extern "C" magma_err_t
magma_clarfb2_gpu( magma_int_t m, magma_int_t n, magma_int_t k,
const magmaFloatComplex_ptr dV, size_t dV_offset, magma_int_t ldv,
const magmaFloatComplex_ptr dT, size_t dT_offset, magma_int_t ldt,
magmaFloatComplex_ptr dC, size_t dC_offset, magma_int_t ldc,
magmaFloatComplex_ptr dwork, size_t dwork_offset, magma_int_t ldwork,
magma_queue_t queue )
{
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
if (m <= 0 || n <= 0)
return MAGMA_SUCCESS;
// W = C^H V
// magma_cgemm( MagmaConjTrans, MagmaNoTrans,
magmablas_cgemm_reduce(
n, k, m,
c_one, dC, dC_offset, ldc,
dV, dV_offset, ldv,
c_zero, dwork, dwork_offset, ldwork, queue);
// W = W T^H = C^H V T^H
magma_ctrmm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
n, k,
c_one, dT, dT_offset, ldt,
dwork, dwork_offset, ldwork,
queue);
// C = C - V W^H = C - V T V^H C = (I - V T V^H) C = H C
magma_cgemm( MagmaNoTrans, MagmaConjTrans,
m, n, k,
c_neg_one, dV, dV_offset, ldv,
dwork, dwork_offset, ldwork,
c_one, dC, dC_offset, ldc, queue );
}
void gmm_magma(const Tensor_core<complex<float>,2>& A, const Tensor_core<complex<float>,2>& B, Tensor_core<complex<float>,2>& C,
char TRANSA, char TRANSB, complex<float> alpha, complex<float> beta)
{
int AL0 = A.rank(0); int AL1 = A.rank(1);
int BL0 = B.rank(0); int BL1 = B.rank(1);
int CL0 = C.rank(0); int CL1 = C.rank(1);
magma_int_t M, N, K, LDA, LDB, LDC;
magma_trans_t transA=magma_trans_const(TRANSA), transB=magma_trans_const(TRANSB);
magmaFloatComplex_ptr d_A, d_B, d_C;
//Set LDA, LDB, and LDC, round up to multiple of 32 for best GPU performance
LDA = ((AL0+31)/32)*32; LDB = ((BL0+31)/32)*32; LDC = ((CL0+31)/32)*32;
// Allocate memory for the matrices on GPU
magma_cmalloc(&d_A, LDA*AL1 );
magma_cmalloc(&d_B, LDB*BL1 );
magma_cmalloc(&d_C, LDC*CL1 );
// Copy data from host (CPU) to device (GPU)
magma_csetmatrix( AL0, AL1, (magmaFloatComplex* ) A.data(), AL0, d_A, LDA );
magma_csetmatrix( BL0, BL1, (magmaFloatComplex* ) B.data(), BL0, d_B, LDB );
if( abs(beta)>1e-32 ) magma_csetmatrix( CL0, CL1, (magmaFloatComplex* ) C.data(), CL0, d_C, LDC );
//Call magma_sgemm
M=( TRANSA=='N' || TRANSA=='n' ) ? AL0:AL1;
K=( TRANSA=='N' || TRANSA=='n' ) ? AL1:AL0;
N=( TRANSB=='N' || TRANSB=='n' ) ? BL1:BL0;
magma_cgemm(transA, transB, M, N, K, _cast_C(alpha), d_A, LDA, d_B, LDB, _cast_C(beta),d_C, LDC);
// Copy solution from device (GPU) to host (CPU)
magma_cgetmatrix(CL0, CL1, d_C, LDC, (magmaFloatComplex* ) C.data(), CL0);
// Free memory on GPU
magma_free(d_A); magma_free(d_B); magma_free(d_C);
}
/**
Purpose
-------
CGERFS improve the computed solution to a system of linear
equations.
The iterative refinement process is stopped if
ITER > ITERMAX
or for all the RHS we have:
RNRM < SQRT(n)*XNRM*ANRM*EPS*BWDMAX
where
o ITER is the number of the current iteration in the iterative
refinement process
o RNRM is the infinity-norm of the residual
o XNRM is the infinity-norm of the solution
o ANRM is the infinity-operator-norm of the matrix A
o EPS is the machine epsilon returned by SLAMCH('Epsilon')
The value ITERMAX and BWDMAX are fixed to 30 and 1.0D+00 respectively.
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 number of linear equations, i.e., 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 array on the GPU, dimension (ldda,N)
the N-by-N coefficient matrix A.
@param[in]
ldda INTEGER
The leading dimension of the array dA. ldda >= max(1,N).
@param[in]
dB COMPLEX array on the GPU, dimension (lddb,NRHS)
The N-by-NRHS right hand side matrix B.
@param[in]
lddb INTEGER
The leading dimension of the array dB. lddb >= max(1,N).
@param[in, out]
dX COMPLEX array on the GPU, dimension (lddx,NRHS)
On entry, the solution matrix X, as computed by
CGETRS_NOPIV. On exit, the improved solution matrix X.
@param[in]
lddx INTEGER
The leading dimension of the array dX. lddx >= max(1,N).
@param
dworkd (workspace) COMPLEX array on the GPU, dimension (N*NRHS)
This array is used to hold the residual vectors.
@param
dAF COMPLEX array on the GPU, dimension (ldda,n)
The factors L and U from the factorization A = L*U
as computed by CGETRF_NOPIV.
@param[out]
iter INTEGER
- < 0: iterative refinement has failed, real
factorization has been performed
+ -1 : the routine fell back to full precision for
implementation- or machine-specific reasons
+ -2 : narrowing the precision induced an overflow,
the routine fell back to full precision
+ -3 : failure of SGETRF
+ -31: stop the iterative refinement after the 30th iteration
- > 0: iterative refinement has been successfully used.
Returns the number of iterations
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if info = -i, the i-th argument had an illegal value
- > 0: if info = i, U(i,i) computed in REAL is
exactly zero. The factorization has been completed,
but the factor U is exactly singular, so the solution
could not be computed.
@ingroup magma_cgesv_driver
********************************************************************/
extern "C" magma_int_t
magma_cgerfs_nopiv_gpu(
magma_trans_t trans, magma_int_t n, magma_int_t nrhs,
magmaFloatComplex_ptr dA, magma_int_t ldda,
magmaFloatComplex_ptr dB, magma_int_t lddb,
magmaFloatComplex_ptr dX, magma_int_t lddx,
magmaFloatComplex_ptr dworkd, magmaFloatComplex_ptr dAF,
magma_int_t *iter,
magma_int_t *info)
{
#define dB(i,j) (dB + (i) + (j)*lddb)
#define dX(i,j) (dX + (i) + (j)*lddx)
#define dR(i,j) (dR + (i) + (j)*lddr)
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_int_t ione = 1;
magmaFloatComplex_ptr dR;
magmaFloatComplex Xnrmv, Rnrmv;
float Anrm, Xnrm, Rnrm, cte, eps;
magma_int_t i, j, iiter, lddsa, lddr;
/* Check arguments */
*iter = 0;
*info = 0;
if ( n < 0 )
*info = -1;
else if ( nrhs < 0 )
*info = -2;
else if ( ldda < max(1,n))
*info = -4;
else if ( lddb < max(1,n))
*info = -8;
else if ( lddx < max(1,n))
*info = -10;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if ( n == 0 || nrhs == 0 )
return *info;
lddsa = n;
lddr = n;
dR = dworkd;
eps = lapackf77_slamch("Epsilon");
Anrm = magmablas_clange(MagmaInfNorm, n, n, dA, ldda, (float*)dworkd );
cte = Anrm * eps * pow( (float)n, (float)0.5 ) * BWDMAX;
// residual dR = dB - dA*dX in real
magmablas_clacpy( MagmaUpperLower, n, nrhs, dB, lddb, dR, lddr );
if ( nrhs == 1 ) {
magma_cgemv( trans, n, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_cgemm( trans, MagmaNoTrans, n, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
// TODO: use MAGMA_C_ABS( dX(i,j) ) instead of clange?
for( j=0; j < nrhs; j++ ) {
i = magma_icamax( n, dX(0,j), 1) - 1;
magma_cgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_clange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_icamax ( n, dR(0,j), 1 ) - 1;
magma_cgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_clange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
// printf("Rnrm : %e, Xnrm*cte : %e\n", Rnrm, Xnrm*cte);
if ( Rnrm > Xnrm*cte ) {
goto REFINEMENT;
}
}
*iter = 0;
return *info;
REFINEMENT:
for( iiter=1; iiter < ITERMAX; ) {
*info = 0;
// solve dAF*dX = dR
// it's okay that dR is used for both dB input and dX output.
magma_cgetrs_nopiv_gpu( trans, n, nrhs, dAF, lddsa, dR, lddr, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// Add correction and setup residual
// dX += dR --and--
// dR = dB
// This saves going through dR a second time (if done with one more kernel).
// -- not really: first time is read, second time is write.
for( j=0; j < nrhs; j++ ) {
magmablas_caxpycp2( n, dR(0,j), dX(0,j), dB(0,j) );
}
// residual dR = dB - dA*dX in real
if ( nrhs == 1 ) {
magma_cgemv( trans, n, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_cgemm( trans, MagmaNoTrans, n, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
/* Check whether the nrhs normwise backward errors satisfy the
* stopping criterion. If yes, set ITER=IITER > 0 and return. */
for( j=0; j < nrhs; j++ ) {
i = magma_icamax( n, dX(0,j), 1) - 1;
magma_cgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_clange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_icamax ( n, dR(0,j), 1 ) - 1;
magma_cgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_clange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto L20;
}
}
/* If we are here, the nrhs normwise backward errors satisfy
* the stopping criterion, we are good to exit. */
*iter = iiter;
return *info;
L20:
iiter++;
}
/* If we are at this place of the code, this is because we have
* performed ITER=ITERMAX iterations and never satisified the
* stopping criterion. Set up the ITER flag accordingly. */
*iter = -ITERMAX - 1;
FALLBACK:
/* Iterative refinement failed to converge to a
* satisfactory solution. */
return *info;
}
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, t1, t2;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t ione = 1;
magma_trans_t trans[] = { MagmaNoTrans, MagmaConjTrans, MagmaTrans };
magma_uplo_t uplo [] = { MagmaLower, MagmaUpper };
magma_diag_t diag [] = { MagmaUnit, MagmaNonUnit };
magma_side_t side [] = { MagmaLeft, MagmaRight };
magmaFloatComplex *A, *B, *C, *C2, *LU;
magmaFloatComplex_ptr dA, dB, dC1, dC2;
magmaFloatComplex alpha = MAGMA_C_MAKE( 0.5, 0.1 );
magmaFloatComplex beta = MAGMA_C_MAKE( 0.7, 0.2 );
float dalpha = 0.6;
float dbeta = 0.8;
float work[1], error, total_error;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t m, n, k, size, maxn, ld, info;
magma_int_t *piv;
magma_int_t err;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf( "Compares magma wrapper function to cublas function; all diffs should be exactly 0.\n\n" );
total_error = 0.;
for( int itest = 0; itest < opts.ntest; ++itest ) {
m = opts.msize[itest];
n = opts.nsize[itest];
k = opts.ksize[itest];
printf("=========================================================================\n");
printf( "m=%d, n=%d, k=%d\n", (int) m, (int) n, (int) k );
// allocate matrices
// over-allocate so they can be any combination of {m,n,k} x {m,n,k}.
maxn = max( max( m, n ), k );
ld = max( 1, maxn );
size = ld*maxn;
err = magma_malloc_cpu( (void**) &piv, maxn*sizeof(magma_int_t) ); assert( err == 0 );
err = magma_cmalloc_pinned( &A, size ); assert( err == 0 );
err = magma_cmalloc_pinned( &B, size ); assert( err == 0 );
err = magma_cmalloc_pinned( &C, size ); assert( err == 0 );
err = magma_cmalloc_pinned( &C2, size ); assert( err == 0 );
err = magma_cmalloc_pinned( &LU, size ); assert( err == 0 );
err = magma_cmalloc( &dA, size ); assert( err == 0 );
err = magma_cmalloc( &dB, size ); assert( err == 0 );
err = magma_cmalloc( &dC1, size ); assert( err == 0 );
err = magma_cmalloc( &dC2, size ); assert( err == 0 );
// initialize matrices
size = maxn*maxn;
lapackf77_clarnv( &ione, ISEED, &size, A );
lapackf77_clarnv( &ione, ISEED, &size, B );
lapackf77_clarnv( &ione, ISEED, &size, C );
printf( "========== Level 1 BLAS ==========\n" );
// ----- test CSWAP
// swap columns 2 and 3 of dA, then copy to C2 and compare with A
if ( n >= 3 ) {
magma_csetmatrix( m, n, A, ld, dA, ld );
magma_csetmatrix( m, n, A, ld, dB, ld );
magma_cswap( m, dA(0,1), 1, dA(0,2), 1 );
magma_cswap( m, dB(0,1), 1, dB(0,2), 1 );
// check results, storing diff between magma and cuda calls in C2
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dA, 1, dB, 1 );
magma_cgetmatrix( m, n, dB, ld, C2, ld );
error = lapackf77_clange( "F", &m, &k, C2, &ld, work );
total_error += error;
printf( "cswap diff %.2g\n", error );
}
else {
printf( "cswap skipped for n < 3\n" );
}
// ----- test ICAMAX
// get argmax of column of A
magma_csetmatrix( m, k, A, ld, dA, ld );
error = 0;
for( int j = 0; j < k; ++j ) {
magma_int_t i1 = magma_icamax( m, dA(0,j), 1 );
int i2; // NOT magma_int_t, for cublas
cublasIcamax( opts.handle, m, dA(0,j), 1, &i2 );
// todo need sync here?
assert( i1 == i2 );
error += abs( i1 - i2 );
}
total_error += error;
gflops = (float)m * k / 1e9;
printf( "icamax diff %.2g\n", error );
printf( "\n" );
printf( "========== Level 2 BLAS ==========\n" );
// ----- test CGEMV
// c = alpha*A*b + beta*c, with A m*n; b,c m or n-vectors
// try no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
magma_csetmatrix( m, n, A, ld, dA, ld );
magma_csetvector( maxn, B, 1, dB, 1 );
magma_csetvector( maxn, C, 1, dC1, 1 );
magma_csetvector( maxn, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_cgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasCgemv( opts.handle, cublas_trans_const(trans[ia]),
m, n, &alpha, dA, ld, dB, 1, &beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
size = (trans[ia] == MagmaNoTrans ? m : n);
cublasCaxpy( opts.handle, size, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetvector( size, dC2, 1, C2, 1 );
error = lapackf77_clange( "F", &size, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_CGEMV( m, n ) / 1e9;
printf( "cgemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_trans_const(trans[ia]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test CHEMV
// c = alpha*A*b + beta*c, with A m*m symmetric; b,c m-vectors
// try upper/lower
for( int iu = 0; iu < 2; ++iu ) {
magma_csetmatrix( m, m, A, ld, dA, ld );
magma_csetvector( m, B, 1, dB, 1 );
magma_csetvector( m, C, 1, dC1, 1 );
magma_csetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_chemv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasChemv( opts.handle, cublas_uplo_const(uplo[iu]),
m, &alpha, dA, ld, dB, 1, &beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasCaxpy( opts.handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_clange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_CHEMV( m ) / 1e9;
printf( "chemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test CTRSV
// solve A*c = c, with A m*m triangular; c m-vector
// try upper/lower, no-trans/trans, unit/non-unit diag
// Factor A into LU to get well-conditioned triangles, else solve yields garbage.
// Still can give garbage if solves aren't consistent with LU factors,
// e.g., using unit diag for U, so copy lower triangle to upper triangle.
// Also used for trsm later.
lapackf77_clacpy( "Full", &maxn, &maxn, A, &ld, LU, &ld );
lapackf77_cgetrf( &maxn, &maxn, LU, &ld, piv, &info );
for( int j = 0; j < maxn; ++j ) {
for( int i = 0; i < j; ++i ) {
*LU(i,j) = *LU(j,i);
}
}
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
magma_csetmatrix( m, m, LU, ld, dA, ld );
magma_csetvector( m, C, 1, dC1, 1 );
magma_csetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_ctrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasCtrsv( opts.handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
cublas_diag_const(diag[id]), m, dA, ld, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasCaxpy( opts.handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_clange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_CTRSM( MagmaLeft, m, 1 ) / 1e9;
printf( "ctrsv( %c, %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]), lapacke_diag_const(diag[id]),
error, gflops/t1, gflops/t2 );
}}}
printf( "\n" );
printf( "========== Level 3 BLAS ==========\n" );
// ----- test CGEMM
// C = alpha*A*B + beta*C, with A m*k or k*m; B k*n or n*k; C m*n
// try combinations of no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
for( int ib = 0; ib < 3; ++ib ) {
bool nta = (trans[ia] == MagmaNoTrans);
bool ntb = (trans[ib] == MagmaNoTrans);
magma_csetmatrix( (nta ? m : k), (nta ? m : k), A, ld, dA, ld );
magma_csetmatrix( (ntb ? k : n), (ntb ? n : k), B, ld, dB, ld );
magma_csetmatrix( m, n, C, ld, dC1, ld );
magma_csetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_cgemm( trans[ia], trans[ib], m, n, k, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasCgemm( opts.handle, cublas_trans_const(trans[ia]), cublas_trans_const(trans[ib]),
m, n, k, &alpha, dA, ld, dB, ld, &beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CGEMM( m, n, k ) / 1e9;
printf( "cgemm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_trans_const(trans[ia]), lapacke_trans_const(trans[ib]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test CHEMM
// C = alpha*A*B + beta*C (left) with A m*m symmetric; B,C m*n; or
// C = alpha*B*A + beta*C (right) with A n*n symmetric; B,C m*n
// try left/right, upper/lower
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
magma_csetmatrix( m, m, A, ld, dA, ld );
magma_csetmatrix( m, n, B, ld, dB, ld );
magma_csetmatrix( m, n, C, ld, dC1, ld );
magma_csetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_chemm( side[is], uplo[iu], m, n, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasChemm( opts.handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
m, n, &alpha, dA, ld, dB, ld, &beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CHEMM( side[is], m, n ) / 1e9;
printf( "chemm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_side_const(side[is]), lapacke_uplo_const(uplo[iu]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test CHERK
// C = alpha*A*A^H + beta*C (no-trans) with A m*k and C m*m symmetric; or
// C = alpha*A^H*A + beta*C (trans) with A k*m and C m*m symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
magma_csetmatrix( n, k, A, ld, dA, ld );
magma_csetmatrix( n, n, C, ld, dC1, ld );
magma_csetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_cherk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasCherk( opts.handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
n, k, &dalpha, dA, ld, &dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CHERK( k, n ) / 1e9;
printf( "cherk( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test CHER2K
// C = alpha*A*B^H + ^alpha*B*A^H + beta*C (no-trans) with A,B n*k; C n*n symmetric; or
// C = alpha*A^H*B + ^alpha*B^H*A + beta*C (trans) with A,B k*n; C n*n symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
bool nt = (trans[it] == MagmaNoTrans);
magma_csetmatrix( (nt ? n : k), (nt ? n : k), A, ld, dA, ld );
magma_csetmatrix( n, n, C, ld, dC1, ld );
magma_csetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_cher2k( uplo[iu], trans[it], n, k, alpha, dA, ld, dB, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasCher2k( opts.handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
n, k, &alpha, dA, ld, dB, ld, &dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CHER2K( k, n ) / 1e9;
printf( "cher2k( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test CTRMM
// C = alpha*A*C (left) with A m*m triangular; C m*n; or
// C = alpha*C*A (right) with A n*n triangular; C m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == MagmaLeft);
magma_csetmatrix( (left ? m : n), (left ? m : n), A, ld, dA, ld );
magma_csetmatrix( m, n, C, ld, dC1, ld );
magma_csetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ctrmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
// note cublas does trmm out-of-place (i.e., adds output matrix C),
// but allows C=B to do in-place.
t2 = magma_sync_wtime( 0 );
cublasCtrmm( opts.handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
cublas_trans_const(trans[it]), cublas_diag_const(diag[id]),
m, n, &alpha, dA, ld, dC2, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CTRMM( side[is], m, n ) / 1e9;
printf( "ctrmm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// ----- test CTRSM
// solve A*X = alpha*B (left) with A m*m triangular; B m*n; or
// solve X*A = alpha*B (right) with A n*n triangular; B m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == MagmaLeft);
magma_csetmatrix( (left ? m : n), (left ? m : n), LU, ld, dA, ld );
magma_csetmatrix( m, n, C, ld, dC1, ld );
magma_csetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ctrsm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasCtrsm( opts.handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
cublas_trans_const(trans[it]), cublas_diag_const(diag[id]),
m, n, &alpha, dA, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasCaxpy( opts.handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_cgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_clange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_CTRSM( side[is], m, n ) / 1e9;
printf( "ctrsm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// cleanup
magma_free_cpu( piv );
magma_free_pinned( A );
magma_free_pinned( B );
magma_free_pinned( C );
magma_free_pinned( C2 );
magma_free_pinned( LU );
magma_free( dA );
magma_free( dB );
magma_free( dC1 );
magma_free( dC2 );
fflush( stdout );
}
if ( total_error != 0. ) {
printf( "total error %.2g -- ought to be 0 -- some test failed (see above).\n",
total_error );
}
else {
printf( "all tests passed\n" );
}
TESTING_FINALIZE();
int status = (total_error != 0.);
return status;
}
extern "C" magma_int_t
magma_cpotrf2_msub(
magma_int_t num_subs, magma_int_t num_gpus, magma_uplo_t uplo, magma_int_t m, magma_int_t n,
magma_int_t off_i, magma_int_t off_j, magma_int_t nb,
magmaFloatComplex_ptr *d_lA, size_t d_lA_offset, magma_int_t ldda,
magmaFloatComplex_ptr *d_lP, magma_int_t lddp,
magmaFloatComplex *a, magma_int_t lda, magma_int_t h,
magma_queue_t *queues,
magma_int_t *info )
{
/* -- clMAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date November 2014
Purpose
=======
CPOTRF computes the Cholesky factorization of a complex Hermitian
positive definite matrix dA.
The factorization has the form
dA = U**H * U, if UPLO = 'U', or
dA = L * L**H, if UPLO = 'L',
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': Upper triangle of dA is stored;
= 'L': Lower triangle of dA is stored.
N (input) INTEGER
The order of the matrix dA. N >= 0.
dA (input/output) COMPLEX array on the GPU, dimension (LDDA,N)
On entry, the Hermitian matrix dA. If UPLO = 'U', the leading
N-by-N upper triangular part of dA contains the upper
triangular part of the matrix dA, and the strictly lower
triangular part of dA is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of dA contains the lower
triangular part of the matrix dA, and the strictly upper
triangular part of dA is not referenced.
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization dA = U**H * U or dA = L * L**H.
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
===================================================================== */
int tot_subs = num_subs*num_gpus;
magma_int_t j, jb, nb0, nb2, dd, d, id, j_local, j_local2;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
float d_one = 1.0;
float d_neg_one = -1.0;
magmaFloatComplex_ptr dlpanel;
size_t dlpanel_offset;
magma_int_t n_local[MagmaMaxSubs * MagmaMaxGPUs], ldpanel;
// initialize trace
trace_init(1, num_gpus, 2, queues);
*info = 0;
if ( (uplo != MagmaUpper) && (uplo != MagmaLower) ) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if ((uplo != MagmaUpper) && tot_subs*ldda < max(1,n)) {
*info = -4;
} else if ((uplo == MagmaUpper) && ldda < max(1,m)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
for (d=0; d<tot_subs; d++) {
/* local-n and local-ld */
if (uplo == MagmaUpper) {
n_local[d] = ((n/nb)/tot_subs)*nb;
if (d < (n/nb)%tot_subs)
n_local[d] += nb;
else if (d == (n/nb)%tot_subs)
n_local[d] += n%nb;
} else {
n_local[d] = ((m/nb)/tot_subs)*nb;
if (d < (m/nb)%tot_subs)
n_local[d] += nb;
else if (d == (m/nb)%tot_subs)
n_local[d] += m%nb;
}
}
/* Use blocked code. */
if (uplo == MagmaUpper) {
/* ---------------------------------------------- */
/* Upper-triangular case */
/* > Compute the Cholesky factorization A = U'*U. */
/* ---------------------------------------------- */
for (j=0; j<m; j+=nb) {
/* Set the GPU number that holds the current panel */
id = (j/nb)%tot_subs;
/* Set the local index where the current panel is */
j_local = j/(nb*tot_subs);
jb = min(nb, (m-j));
if (j > 0) {
// Wait for the column on CPU
magma_queue_sync(queues[2*(id%num_gpus)]);
/* broadcast off-diagonal column to all gpus */
d = (j/nb+1)%num_gpus;
for (dd=0; dd<num_gpus; dd++) {
if (d != id%num_gpus) {
magma_csetmatrix_async( j, jb,
Aup(0,j), lda,
dlP(d,jb,0,id%num_gpus), lddp,
queues[2*d],
trace_gpu_event(d, 0, "set", "set-col") );
}
d = (d+1)%num_gpus;
}
/* Update the current diagonal block */
trace_gpu_start(id%num_gpus, 1, "herk", "herk");
magma_cherk(MagmaUpper, MagmaConjTrans, jb, j,
d_neg_one, dlA(id, 0, nb*j_local), ldda,
d_one, dlA(id, j, nb*j_local), ldda,
queues[2*(id%num_gpus)+1]);
magma_queue_sync(queues[2*(id%num_gpus)+1]); // Wait for syrk
}
/* Send the diagonal to cpu */
magma_cgetmatrix_async( jb, jb,
dlA(id, j, nb*j_local), ldda,
Aup(j,j), lda,
queues[2*(id%num_gpus)],
trace_gpu_event(id%num_gpus, 0, "get", "get-diag") );
if (j > 0) {
/* Compute the local block column of the panel. */
d = (j/nb+1)%tot_subs;
for (dd=0; dd<tot_subs; dd++) {
j_local2 = j_local+1;
if (d > id) j_local2 --;
nb0 = nb*j_local2;
if (n_local[d] > nb0) {
if (d%num_gpus != id%num_gpus) {
dlpanel = d_lP[d%num_gpus];
dlpanel_offset = dlP_offset(jb, 0, id%num_gpus);
ldpanel = lddp;
/* Wait for the offdiagonal column */
if (dd < num_gpus) magma_queue_sync(queues[2*(d%num_gpus)]);
} else {
dlpanel = d_lA[id];
dlpanel_offset = dlA_offset(0, nb*j_local);
ldpanel = ldda;
}
/* update the panel */
trace_gpu_start(d%num_gpus, 1, "gemm", "gemm");
magma_cgemm(MagmaConjTrans, MagmaNoTrans,
jb, n_local[d]-nb0, j,
c_neg_one, dlpanel, dlpanel_offset, ldpanel,
dlA(d, 0, nb0), ldda,
c_one, dlA(d, j, nb0), ldda,
queues[2*(d%num_gpus)+1]);
}
d = (d+1)%tot_subs;
}
}
/* factor the diagonal */
magma_queue_sync( queues[2*(id%num_gpus)] ); // wait for the diagonal
trace_cpu_start(0, "potrf", "potrf");
lapackf77_cpotrf(MagmaUpperStr, &jb, Aup(j,j), &lda, info);
trace_cpu_end(0);
if (*info != 0) {
*info = *info + j;
break;
}
/* send the diagonal to gpus */
if ((j+jb) < n) {
d = (j/nb+1)%num_gpus;
for (dd=0; dd<num_gpus; dd++) {
if (d == id%num_gpus) {
dlpanel = d_lA[id];
dlpanel_offset = dlA_offset(j, nb*j_local);
ldpanel = ldda;
} else {
dlpanel = d_lP[d];
dlpanel_offset = dlP_offset(0, 0, id%num_gpus);
ldpanel = lddp;
}
magma_csetmatrix_async( jb, jb,
Aup(j,j), lda,
dlpanel, dlpanel_offset, ldpanel,
queues[2*d],
trace_gpu_event(d, 0, "set", "set-diag"));
d = (d+1)%num_gpus;
}
} else {
magma_csetmatrix_async( jb, jb,
Aup(j,j), lda,
dlA(id, j, nb*j_local), ldda,
queues[2*(id%num_gpus)],
trace_gpu_event(id%num_gpus, 0, "set", "set-diag") );
}
/* panel-factorize the off-diagonal */
if ((j+jb) < n) {
d = (j/nb+1)%tot_subs;
for (dd=0; dd<tot_subs; dd++) {
/* next column */
j_local2 = j_local+1;
if (d > id) j_local2--;
if (d%num_gpus == id%num_gpus) {
dlpanel = d_lA[id];
dlpanel_offset = dlA_offset(j, nb*j_local);
ldpanel = ldda;
} else {
dlpanel = d_lP[d%num_gpus];
dlpanel_offset = dlP_offset(0, 0, id%num_gpus);
ldpanel = lddp;
}
nb2 = n_local[d]-nb*j_local2;
nb0 = min(nb, nb2);
if (dd < num_gpus) magma_queue_sync( queues[2*(d%num_gpus)] ); // wait for the diagonal
if (j+jb < m && d == (j/nb+1)%tot_subs) {
/* owns the next column, look-ahead the column */
trace_gpu_start(d%num_gpus, 1, "trsm", "trsm");
magma_ctrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb0, c_one,
dlpanel, dlpanel_offset, ldpanel,
dlA(d, j, nb*j_local2), ldda,
queues[2*(d%num_gpus)+1] );
/* send the column to cpu */
magma_queue_sync(queues[2*(d%num_gpus)+1]); // wait for lookahead
magma_cgetmatrix_async( (j+jb), nb0,
dlA(d, 0, nb*j_local2), ldda,
Aup(0,j+jb), lda,
queues[2*(d%num_gpus)],
trace_gpu_event(d%num_gpus, 0, "get", "get-col") );
/* update the remaining blocks */
nb2 = nb2 - nb0;
trace_gpu_start(d%num_gpus, 1, "trsm", "trsm");
magma_ctrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, dlpanel_offset, ldpanel,
dlA(d, j, nb*j_local2+nb0), ldda,
queues[2*(d%num_gpus)+1] );
} else if (nb2 > 0) {
/* update the entire trailing matrix */
trace_gpu_start(d%num_gpus, 1, "trsm", "trsm");
magma_ctrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, dlpanel_offset, ldpanel,
dlA(d, j, nb*j_local2), ldda,
queues[2*(d%num_gpus)+1] );
}
d = (d+1)%tot_subs;
}
}
}
} else {
/* -------------------------------------------- */
/* Lower-triangular case */
/* Compute the Cholesky factorization A = L*L'. */
/* -------------------------------------------- */
for (j=0; j<n; j+=nb) {
/* Set the GPU number that holds the current panel */
id = (j/nb)%tot_subs;
/* Set the local index where the current panel is */
j_local = j/(nb*tot_subs);
jb = min(nb, (n-j));
if (j > 0) {
if (num_gpus > 1) {
// Wait for the row on CPU to broadcast
magma_queue_sync(queues[2*(id%num_gpus)]);
}
/* broadcast off-diagonal row to all the GPUs */
d = (j/nb+1)%num_gpus;
for (dd=0; dd<num_gpus; dd++) {
if (d != id%num_gpus) {
/* send it to GPU-d */
magma_csetmatrix_async( jb, j,
Alo(j,0), lda,
dlPT(d,0,jb,id%num_gpus), nb,
queues[2*d],
trace_gpu_event(d, 0, "set", "set-row") );
}
d = (d+1)%num_gpus;
}
/* Update the current diagonal block */
trace_gpu_start(id%num_gpus, 1, "herk", "herk");
magma_cherk(MagmaLower, MagmaNoTrans, jb, j,
d_neg_one, dlA(id, nb*j_local, 0), ldda,
d_one, dlA(id, nb*j_local, j), ldda,
queues[2*(id%num_gpus)+1]);
magma_queue_sync(queues[2*(id%num_gpus)+1]); // wait for syrk
}
/* send the diagonal to cpu */
magma_cgetmatrix_async( jb, jb,
dlA(id, nb*j_local, j), ldda,
Alo(j,j), lda,
queues[2*(id%num_gpus)],
trace_gpu_event(id%num_gpus, 0, "get", "get") );
/* update the offdiagonal blocks */
if (j > 0) {
/* compute the block-rows of the panel */
d = (j/nb+1)%tot_subs;
for (dd=0; dd<tot_subs; dd++) {
j_local2 = j_local+1;
if (d > id) j_local2 --;
nb0 = nb*j_local2;
if (nb0 < n_local[d]) {
if (d%num_gpus != id%num_gpus) {
dlpanel = d_lP[d%num_gpus];
dlpanel_offset = dlPT_offset(0, jb, id%num_gpus);
ldpanel = nb;
/* Wait for offdiagonal row */
if (dd < num_gpus) magma_queue_sync(queues[2*(d%num_gpus)]);
} else {
dlpanel = d_lA[id];
dlpanel_offset = dlA_offset(nb*j_local, 0);
ldpanel = ldda;
}
/* Update the panel */
trace_gpu_start(d%num_gpus, 1, "gemm", "gemm");
magma_cgemm( MagmaNoTrans, MagmaConjTrans,
n_local[d]-nb0, jb, j,
c_neg_one, dlA(d, nb0, 0), ldda,
dlpanel, dlpanel_offset, ldpanel,
c_one, dlA(d, nb0, j), ldda,
queues[2*(d%num_gpus)+1]);
}
d = (d+1)%tot_subs;
}
}
/* factor the diagonal */
magma_queue_sync( queues[2*(id%num_gpus)] );
trace_cpu_start(0, "potrf", "potrf");
lapackf77_cpotrf(MagmaLowerStr, &jb, Alo(j,j), &lda, info);
trace_cpu_end(0);
if (*info != 0) {
printf( " cpotrf returned %d (id=%d,j=%d,j_local=%d,jb=%d)\n",*info,id,j,j_local,jb );
*info = *info + j;
break;
}
/* send the diagonal to gpus */
if ((j+jb) < m) {
d = (j/nb+1)%num_gpus;
for (dd=0; dd<num_gpus; dd++) {
if (d == id%num_gpus) {
dlpanel = d_lA[id];
dlpanel_offset = dlA_offset(nb*j_local, j);
ldpanel = ldda;
} else {
dlpanel = d_lP[d];
dlpanel_offset = dlPT_offset(0, 0, id%num_gpus);
ldpanel = nb;
}
magma_csetmatrix_async( jb, jb,
Alo(j,j), lda,
dlpanel, dlpanel_offset, ldpanel,
queues[2*d],
trace_gpu_event(d, 0, "set", "set-diag") );
d = (d+1)%num_gpus;
}
} else {
magma_csetmatrix_async( jb, jb,
Alo(j,j), lda,
dlA(id, nb*j_local, j), ldda,
queues[2*(id%num_gpus)],
trace_gpu_event(id%num_gpus, 0, "set", "set-diag") );
}
/* factorize off-diagonal blocks */
if ((j+jb) < m) {
d = (j/nb+1)%tot_subs;
for (dd=0; dd<tot_subs; dd++) {
/* next column */
j_local2 = j_local+1;
if (d > id) j_local2--;
if (d%num_gpus == id%num_gpus) {
dlpanel = d_lA[id];
dlpanel_offset = dlA_offset(nb*j_local, j);
ldpanel = ldda;
} else {
dlpanel = d_lP[d%num_gpus];
dlpanel_offset = dlPT_offset(0, 0, id%num_gpus);
ldpanel = nb;
}
nb2 = n_local[d] - j_local2*nb;
nb0 = min(nb, nb2 );
// wait for the diagonal
if (dd < num_gpus) magma_queue_sync(queues[2*(d%num_gpus)]);
if (j+jb < n && d == (j/nb+1)%tot_subs) {
/* owns the next column, look-ahead the column */
trace_gpu_start(d%num_gpus, 1, "trsm", "trsm");
magma_ctrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb0, jb, c_one,
dlpanel, dlpanel_offset, ldpanel,
dlA(d, nb*j_local2, j), ldda,
queues[2*(d%num_gpus)+1]);
/* send the column to cpu */
magma_queue_sync( queues[2*(d%num_gpus)+1] ); // wait for lookahead
magma_cgetmatrix_async( nb0, j+jb,
dlA(d, nb*j_local2, 0), ldda,
Alo(j+jb,0), lda,
queues[2*(d%num_gpus)],
trace_gpu_event(d%num_gpus, 0, "get", "get") );
/* update the remaining blocks */
nb2 = nb2 - nb0;
trace_gpu_start(d%num_gpus, 1, "trsm", "trsm");
magma_ctrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, dlpanel_offset, ldpanel,
dlA(d, nb*j_local2+nb0, j), ldda,
queues[2*(d%num_gpus)+1]);
} else if (nb2 > 0) {
/* update the entire trailing matrix */
trace_gpu_start(d%num_gpus, 1, "trsm", "trsm");
magma_ctrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, dlpanel_offset, ldpanel,
dlA(d, nb*j_local2, j), ldda,
queues[2*(d%num_gpus)+1]);
}
d = (d+1)%tot_subs;
}
}
}
} /* end of else not upper */
/* clean up */
for( d=0; d<num_gpus; d++ ) {
magma_queue_sync( queues[2*d] );
magma_queue_sync( queues[2*d+1] );
}
trace_finalize("cpotrf_msub.svg", "trace.css");
return *info;
} /* magma_cpotrf2_msub */
/**
Purpose
-------
CSYTRF_nopiv_gpu computes the LDLt factorization of a complex symmetric
matrix A.
The factorization has the form
A = U^T * D * U, if UPLO = MagmaUpper, or
A = L * D * L^T, if UPLO = MagmaLower,
where U is an upper triangular matrix, L is lower triangular, and
D is a diagonal matrix.
This is the block version of the algorithm, calling Level 3 BLAS.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in,out]
dA COMPLEX array on the GPU, dimension (LDDA,N)
On entry, the symmetric matrix A. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
\n
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization A = U^H D U or A = L D L^H.
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using cudaMallocHost.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
if INFO = -6, the GPU memory allocation failed
- > 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
@ingroup magma_csysv_comp
******************************************************************* */
extern "C" magma_int_t
magma_csytrf_nopiv_gpu(
magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex_ptr dA, magma_int_t ldda,
magma_int_t *info)
{
#define A(i, j) (A)
#define dA(i, j) (dA +(j)*ldda + (i))
#define dW(i, j) (dW +(j)*ldda + (i))
#define dWt(i, j) (dW +(j)*nb + (i))
/* Constants */
const magmaFloatComplex c_one = MAGMA_C_ONE;
const magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
/* Local variables */
bool upper = (uplo == MagmaUpper);
magma_int_t j, k, jb, nb, ib, iinfo;
*info = 0;
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (ldda < max(1,n)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ( n == 0 )
return *info;
nb = magma_get_chetrf_nopiv_nb(n);
ib = min(32, nb); // inner-block for diagonal factorization
magma_queue_t queues[2];
magma_event_t event;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
magma_event_create( &event );
trace_init( 1, 1, 2, queues );
// CPU workspace
magmaFloatComplex *A;
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &A, nb*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
// GPU workspace
magmaFloatComplex_ptr dW;
if (MAGMA_SUCCESS != magma_cmalloc( &dW, (1+nb)*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
/* Use hybrid blocked code. */
if (upper) {
//=========================================================
// Compute the LDLt factorization A = U'*D*U without pivoting.
// main loop
for (j=0; j < n; j += nb) {
jb = min(nb, (n-j));
// copy A(j,j) back to CPU
trace_gpu_start( 0, 0, "get", "get" );
magma_event_sync( event );
magma_cgetmatrix_async( jb, jb, dA(j, j), ldda, A(j,j), nb, queues[1] );
trace_gpu_end( 0, 0 );
// factorize the diagonal block
magma_queue_sync( queues[1] );
trace_cpu_start( 0, "potrf", "potrf" );
magma_csytrf_nopiv_cpu( MagmaUpper, jb, ib, A(j, j), nb, info );
trace_cpu_end( 0 );
if (*info != 0) {
*info = *info + j;
break;
}
// copy A(j,j) back to GPU
trace_gpu_start( 0, 0, "set", "set" );
magma_csetmatrix_async( jb, jb, A(j, j), nb, dA(j, j), ldda, queues[0] );
trace_gpu_end( 0, 0 );
if ( (j+jb) < n) {
// compute the off-diagonal blocks of current block column
trace_gpu_start( 0, 0, "trsm", "trsm" );
magma_ctrsm( MagmaLeft, MagmaUpper, MagmaTrans, MagmaUnit,
jb, (n-j-jb),
c_one, dA(j, j), ldda,
dA(j, j+jb), ldda, queues[0] );
magma_ccopymatrix( jb, n-j-jb, dA( j, j+jb ), ldda, dWt( 0, j+jb ), nb, queues[0] );
// update the trailing submatrix with D
magmablas_clascl_diag( MagmaUpper, jb, n-j-jb,
dA(j, j), ldda,
dA(j, j+jb), ldda,
queues[0], &iinfo );
trace_gpu_end( 0, 0 );
// update the trailing submatrix with U and W
trace_gpu_start( 0, 0, "gemm", "gemm" );
for (k=j+jb; k < n; k += nb) {
magma_int_t kb = min(nb,n-k);
magma_cgemm( MagmaTrans, MagmaNoTrans, kb, n-k, jb,
c_neg_one, dWt(0, k), nb,
dA(j, k), ldda,
c_one, dA(k, k), ldda, queues[0] );
if (k == j+jb)
magma_event_record( event, queues[0] );
}
trace_gpu_end( 0, 0 );
}
}
} else {
//=========================================================
// Compute the LDLt factorization A = L*D*L' without pivoting.
// main loop
for (j=0; j < n; j += nb) {
jb = min(nb, (n-j));
// copy A(j,j) back to CPU
trace_gpu_start( 0, 0, "get", "get" );
magma_event_sync( event );
magma_cgetmatrix_async( jb, jb, dA(j, j), ldda, A(j,j), nb, queues[1] );
trace_gpu_end( 0, 0 );
// factorize the diagonal block
magma_queue_sync( queues[1] );
trace_cpu_start( 0, "potrf", "potrf" );
magma_csytrf_nopiv_cpu( MagmaLower, jb, ib, A(j, j), nb, info );
trace_cpu_end( 0 );
if (*info != 0) {
*info = *info + j;
break;
}
// copy A(j,j) back to GPU
trace_gpu_start( 0, 0, "set", "set" );
magma_csetmatrix_async( jb, jb, A(j, j), nb, dA(j, j), ldda, queues[0] );
trace_gpu_end( 0, 0 );
if ( (j+jb) < n) {
// compute the off-diagonal blocks of current block column
trace_gpu_start( 0, 0, "trsm", "trsm" );
magma_ctrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
(n-j-jb), jb,
c_one, dA(j, j), ldda,
dA(j+jb, j), ldda, queues[0] );
magma_ccopymatrix( n-j-jb,jb, dA( j+jb, j ), ldda, dW( j+jb, 0 ), ldda, queues[0] );
// update the trailing submatrix with D
magmablas_clascl_diag(MagmaLower, n-j-jb, jb,
dA(j, j), ldda,
dA(j+jb, j), ldda,
queues[0], &iinfo);
trace_gpu_end( 0, 0 );
// update the trailing submatrix with L and W
trace_gpu_start( 0, 0, "gemm", "gemm" );
for (k=j+jb; k < n; k += nb) {
magma_int_t kb = min(nb,n-k);
magma_cgemm( MagmaNoTrans, MagmaTrans, n-k, kb, jb,
c_neg_one, dA(k, j), ldda,
dW(k, 0), ldda,
c_one, dA(k, k), ldda, queues[0] );
if (k == j+jb)
magma_event_record( event, queues[0] );
}
trace_gpu_end( 0, 0 );
}
}
}
trace_finalize( "chetrf.svg","trace.css" );
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_event_destroy( event );
magma_free( dW );
magma_free_pinned( A );
return *info;
} /* magma_csytrf_nopiv */
/**
Purpose
-------
CGETRF_NOPIV_GPU computes an LU factorization of a general M-by-N
matrix A without any pivoting.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
dA COMPLEX array on the GPU, dimension (LDDA,N).
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
- > 0: if INFO = i, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
@ingroup magma_cgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_cgetrf_nopiv_gpu(
magma_int_t m, magma_int_t n,
magmaFloatComplex_ptr dA, magma_int_t ldda,
magma_int_t *info)
{
#define dA(i,j) (dA + (i)*nb + (j)*nb*ldda)
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t iinfo, nb;
magma_int_t maxm, mindim;
magma_int_t i, rows, s, lddwork;
magmaFloatComplex *work;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
mindim = min(m, n);
nb = magma_get_cgetrf_nb(m);
s = mindim / nb;
if (nb <= 1 || nb >= min(m,n)) {
/* Use CPU code. */
magma_cmalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_cgetmatrix( m, n, dA, ldda, work, m );
magma_cgetrf_nopiv( m, n, work, m, info);
magma_csetmatrix( m, n, work, m, dA, ldda );
magma_free_cpu(work);
}
else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
lddwork = maxm;
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, maxm*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
/* Define user stream if current stream is NULL */
magma_queue_t stream[2];
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
magma_queue_create( &stream[0] );
if (orig_stream == NULL) {
magma_queue_create( &stream[1] );
magmablasSetKernelStream(stream[1]);
}
else {
stream[1] = orig_stream;
}
for( i=0; i < s; i++ ) {
// download i-th panel
magma_queue_sync( stream[1] );
magma_cgetmatrix_async( m-i*nb, nb, dA(i,i), ldda, work, lddwork, stream[0] );
if ( i > 0 ) {
magma_ctrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit,
nb, n - (i+1)*nb,
c_one, dA(i-1,i-1), ldda,
dA(i-1,i+1), ldda );
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
m-i*nb, n-(i+1)*nb, nb,
c_neg_one, dA(i, i-1), ldda, dA(i-1,i+1), ldda,
c_one, dA(i, i+1), ldda );
}
// do the cpu part
rows = m - i*nb;
magma_queue_sync( stream[0] );
magma_cgetrf_nopiv( rows, nb, work, lddwork, &iinfo );
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + i*nb;
// upload i-th panel
magma_csetmatrix_async( m-i*nb, nb, work, lddwork, dA(i, i), ldda, stream[0] );
magma_queue_sync( stream[0] );
// do the small non-parallel computations
if ( s > (i+1) ) {
magma_ctrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dA(i, i ), ldda,
dA(i, i+1), ldda);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
m-(i+1)*nb, nb, nb,
c_neg_one, dA(i+1, i ), ldda, dA(i, i+1), ldda,
c_one, dA(i+1, i+1), ldda );
}
else {
magma_ctrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit,
nb, n-s*nb,
c_one, dA(i, i ), ldda,
dA(i, i+1), ldda);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
m-(i+1)*nb, n-(i+1)*nb, nb,
c_neg_one, dA(i+1, i ), ldda, dA(i, i+1), ldda,
c_one, dA(i+1, i+1), ldda );
}
}
magma_int_t nb0 = min(m - s*nb, n - s*nb);
rows = m - s*nb;
magma_cgetmatrix( rows, nb0, dA(s,s), ldda, work, lddwork );
// make sure that gpu queue is empty
magma_device_sync();
// do the cpu part
magma_cgetrf_nopiv( rows, nb0, work, lddwork, &iinfo );
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + s*nb;
// upload i-th panel
magma_csetmatrix( rows, nb0, work, lddwork, dA(s,s), ldda );
magma_ctrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit,
nb0, n-s*nb-nb0,
c_one, dA(s,s), ldda,
dA(s,s)+nb0, ldda);
magma_free_pinned( work );
magma_queue_destroy( stream[0] );
if (orig_stream == NULL) {
magma_queue_destroy( stream[1] );
}
magmablasSetKernelStream( orig_stream );
}
return *info;
} /* magma_cgetrf_nopiv_gpu */
extern "C" void
magma_cher2k_mgpu(
magma_int_t num_gpus, magma_uplo_t uplo, magma_trans_t trans, magma_int_t nb, magma_int_t n, magma_int_t k,
magmaFloatComplex alpha,
magmaFloatComplex **db, magma_int_t lddb, magma_int_t offset_b,
float beta,
magmaFloatComplex **dc, magma_int_t lddc, magma_int_t offset,
magma_int_t num_streams, magma_queue_t stream[][10])
{
#define dB(id, i, j) (db[(id)]+(j)*lddb + (i)+offset_b)
#define dB1(id, i, j) (db[(id)]+(j)*lddb + (i)+offset_b)+k*lddb
#define dC(id, i, j) (dc[(id)]+(j)*lddc + (i))
magma_int_t i, id, ib, ii, kk, n1;
magmaFloatComplex c_one = MAGMA_C_ONE;
/* diagonal update */
for( i=0; i < n; i += nb ) {
id = ((i+offset)/nb)%num_gpus;
kk = (i/(nb*num_gpus))%num_streams;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
ib = min(nb, n-i);
ii = nb*((i+offset)/(nb*num_gpus));
/* cher2k on diagonal block */
trace_gpu_start( id, kk, "syr2k", "syr2k" );
magma_cher2k(uplo, trans, ib, k,
alpha, dB1(id, i, 0 ), lddb,
dB(id, i, 0 ), lddb,
beta, dC(id, i+offset, ii), lddc);
trace_gpu_end( id, kk );
}
/* off-diagonal update */
if (uplo == MagmaUpper) {
for( i=nb; i < n; i += nb ) {
id = ((i+offset)/nb)%num_gpus;
kk = (i/(nb*num_gpus))%num_streams;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
ib = min(nb, n-i);
ii = nb*((i+offset)/(nb*num_gpus));
magma_cgemm(MagmaNoTrans, MagmaConjTrans, i, ib, k,
alpha, dB1(id, 0, 0 ), lddb,
dB(id, i, 0 ), lddb,
c_one, dC(id, 0, ii), lddc);
}
}
else {
for( i=0; i < n-nb; i += nb ) {
id = ((i+offset)/nb)%num_gpus;
kk = (i/(nb*num_gpus))%num_streams;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
ib = min(nb, n-i);
ii = nb*((i+offset)/(nb*num_gpus));
n1 = n-i-ib;
// cgemm on off-diagonal blocks
trace_gpu_start( id, kk, "gemm_up", "gemm_up" );
magma_cgemm(MagmaNoTrans, MagmaConjTrans, n1, ib, k,
alpha, dB1(id, i+ib, 0 ), lddb,
dB(id, i, 0 ), lddb,
c_one, dC(id, i+offset+ib, ii), lddc);
trace_gpu_end( id, kk );
}
}
if (uplo == MagmaUpper) {
for( i=nb; i < n; i += nb ) {
id = ((i+offset)/nb)%num_gpus;
kk = (i/(nb*num_gpus))%num_streams;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
ib = min(nb, n-i);
ii = nb*((i+offset)/(nb*num_gpus));
magma_cgemm(MagmaNoTrans, MagmaConjTrans, i, ib, k,
alpha, dB( id, 0, 0 ), lddb,
dB1(id, i, 0 ), lddb,
c_one, dC(id, 0, ii), lddc);
}
} else {
for( i=0; i < n-nb; i += nb ) {
id = ((i+offset)/nb)%num_gpus;
kk = (i/(nb*num_gpus))%num_streams;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
ib = min(nb, n-i);
ii = nb*((i+offset)/(nb*num_gpus));
n1 = n-i-ib;
/* cgemm on off-diagonal blocks */
trace_gpu_start( id, kk, "gemm_up", "gemm_up" );
magma_cgemm(MagmaNoTrans, MagmaConjTrans, n1, ib, k,
alpha, dB(id, i+ib, 0 ), lddb,
dB1(id, i, 0 ), lddb,
c_one, dC(id, i+offset+ib, ii), lddc);
trace_gpu_end( id, kk );
}
}
for( id=0; id < num_gpus; id++ ) {
magma_setdevice(id);
for( kk=0; kk < num_streams; kk++ ) magma_queue_sync(stream[id][kk]);
magmablasSetKernelStream(NULL);
}
}
/**
Purpose
-------
CGETRF_m computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges. This version does not
require work space on the GPU passed as input. GPU memory is allocated
in the routine. The matrix may exceed the GPU memory.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
Note: The factorization of big panel is done calling multiple-gpu-interface.
Pivots are applied on GPU within the big panel.
Arguments
---------
@param[in]
ngpu INTEGER
Number of GPUs to use. ngpu > 0.
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
ipiv INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
- > 0: if INFO = i, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
@ingroup magma_cgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_cgetrf_m(
magma_int_t ngpu,
magma_int_t m, magma_int_t n,
magmaFloatComplex *A, magma_int_t lda, magma_int_t *ipiv,
magma_int_t *info)
{
#define A(i,j) (A + (j)*lda + (i))
#define dAT(d,i,j) (dAT[d] + (i)*nb*ldn_local + (j)*nb)
#define dPT(d,i,j) (dPT[d] + (i)*nb*nb + (j)*nb*maxm)
magma_timer_t time=0, time_total=0, time_alloc=0, time_set=0, time_get=0, time_comp=0;
timer_start( time_total );
real_Double_t flops;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *dAT[MagmaMaxGPUs], *dA[MagmaMaxGPUs], *dPT[MagmaMaxGPUs];
magma_int_t iinfo = 0, nb, nbi, maxm, n_local[MagmaMaxGPUs], ldn_local;
magma_int_t N, M, NB, NBk, I, d, ngpu0 = ngpu;
magma_int_t ii, jj, h, offset, ib, rows;
magma_queue_t stream[MagmaMaxGPUs][2];
magma_event_t event[MagmaMaxGPUs][2];
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (lda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
/* initialize nb */
nb = magma_get_cgetrf_nb(m);
maxm = ((m + 31)/32)*32;
/* figure out NB */
size_t freeMem, totalMem;
cudaMemGetInfo( &freeMem, &totalMem );
freeMem /= sizeof(magmaFloatComplex);
/* number of columns in the big panel */
h = 1+(2+ngpu0);
NB = (magma_int_t)(0.8*freeMem/maxm-h*nb);
const char* ngr_nb_char = getenv("MAGMA_NGR_NB");
if ( ngr_nb_char != NULL )
NB = max( nb, min( NB, atoi(ngr_nb_char) ) );
//NB = 5*max(nb,32);
if ( ngpu0 > ceil((float)NB/nb) ) {
ngpu = (int)ceil((float)NB/nb);
h = 1+(2+ngpu);
NB = (magma_int_t)(0.8*freeMem/maxm-h*nb);
} else {
ngpu = ngpu0;
}
if ( ngpu*NB >= n ) {
#ifdef CHECK_CGETRF_OOC
printf( " * still fit in GPU memory.\n" );
#endif
NB = n;
} else {
#ifdef CHECK_CGETRF_OOC
printf( " * don't fit in GPU memory.\n" );
#endif
NB = ngpu*NB;
NB = max( nb, (NB / nb) * nb); /* making sure it's devisable by nb (x64) */
}
#ifdef CHECK_CGETRF_OOC
if ( NB != n ) printf( " * running in out-core mode (n=%d, NB=%d, nb=%d, freeMem=%.2e).\n", n, NB, nb, (float)freeMem );
else printf( " * running in in-core mode (n=%d, NB=%d, nb=%d, freeMem=%.2e).\n", n, NB, nb, (float)freeMem );
#endif
if ( (nb <= 1) || (nb >= min(m,n)) ) {
/* Use CPU code for scalar of one tile. */
lapackf77_cgetrf(&m, &n, A, &lda, ipiv, info);
} else {
/* Use hybrid blocked code. */
/* allocate memory on GPU to store the big panel */
timer_start( time_alloc );
n_local[0] = (NB/nb)/ngpu;
if ( NB%(nb*ngpu) != 0 )
n_local[0]++;
n_local[0] *= nb;
ldn_local = ((n_local[0]+31)/32)*32;
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
if (MAGMA_SUCCESS != magma_cmalloc( &dA[d], (ldn_local+h*nb)*maxm )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
dPT[d] = dA[d] + nb*maxm; /* for storing the previous panel from CPU */
dAT[d] = dA[d] + h*nb*maxm; /* for storing the big panel */
magma_queue_create( &stream[d][0] );
magma_queue_create( &stream[d][1] );
magma_event_create( &event[d][0] );
magma_event_create( &event[d][1] );
}
//magma_setdevice(0);
timer_stop( time_alloc );
for( I=0; I < n; I += NB ) {
M = m;
N = min( NB, n-I ); /* number of columns in this big panel */
//s = min( max(m-I,0), N )/nb; /* number of small block-columns in this big panel */
maxm = ((M + 31)/32)*32;
if ( ngpu0 > ceil((float)N/nb) ) {
ngpu = (int)ceil((float)N/nb);
} else {
ngpu = ngpu0;
}
for( d=0; d < ngpu; d++ ) {
n_local[d] = ((N/nb)/ngpu)*nb;
if (d < (N/nb)%ngpu)
n_local[d] += nb;
else if (d == (N/nb)%ngpu)
n_local[d] += N%nb;
}
ldn_local = ((n_local[0]+31)/32)*32;
/* upload the next big panel into GPU, transpose (A->A'), and pivot it */
timer_start( time );
magmablas_csetmatrix_transpose_mgpu(ngpu, stream, A(0,I), lda,
dAT, ldn_local, dA, maxm, M, N, nb);
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_queue_sync( stream[d][0] );
magma_queue_sync( stream[d][1] );
magmablasSetKernelStream(NULL);
}
time_set += timer_stop( time );
timer_start( time );
/* == --------------------------------------------------------------- == */
/* == loop around the previous big-panels to update the new big-panel == */
for( offset = 0; offset < min(m,I); offset += NB ) {
NBk = min( m-offset, NB );
/* start sending the first tile from the previous big-panels to gpus */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
nbi = min( nb, NBk );
magma_csetmatrix_async( (M-offset), nbi,
A(offset,offset), lda,
dA[d], (maxm-offset), stream[d][0] );
/* make sure the previous update finished */
magmablasSetKernelStream(stream[d][0]);
//magma_queue_sync( stream[d][1] );
magma_queue_wait_event( stream[d][0], event[d][0] );
/* transpose */
magmablas_ctranspose( M-offset, nbi, dA[d], maxm-offset, dPT(d,0,0), nb );
}
/* applying the pivot from the previous big-panel */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magmablas_claswp_q( ldn_local, dAT(d,0,0), ldn_local, offset+1, offset+NBk, ipiv, 1, stream[d][1] );
}
/* == going through each block-column of previous big-panels == */
for( jj=0, ib=offset/nb; jj < NBk; jj += nb, ib++ ) {
ii = offset+jj;
rows = maxm - ii;
nbi = min( nb, NBk-jj );
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
/* wait for a block-column on GPU */
magma_queue_sync( stream[d][0] );
/* start sending next column */
if ( jj+nb < NBk ) {
magma_csetmatrix_async( (M-ii-nb), min(nb,NBk-jj-nb),
A(ii+nb,ii+nb), lda,
dA[d], (rows-nb), stream[d][0] );
/* make sure the previous update finished */
magmablasSetKernelStream(stream[d][0]);
//magma_queue_sync( stream[d][1] );
magma_queue_wait_event( stream[d][0], event[d][(1+jj/nb)%2] );
/* transpose next column */
magmablas_ctranspose( M-ii-nb, nb, dA[d], rows-nb, dPT(d,0,(1+jj/nb)%2), nb );
}
/* update with the block column */
magmablasSetKernelStream(stream[d][1]);
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n_local[d], nbi, c_one, dPT(d,0,(jj/nb)%2), nb, dAT(d,ib,0), ldn_local );
if ( M > ii+nb ) {
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n_local[d], M-(ii+nb), nbi, c_neg_one, dAT(d,ib,0), ldn_local,
dPT(d,1,(jj/nb)%2), nb, c_one, dAT(d,ib+1,0), ldn_local );
}
magma_event_record( event[d][(jj/nb)%2], stream[d][1] );
} /* end of for each block-columns in a big-panel */
}
} /* end of for each previous big-panels */
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_queue_sync( stream[d][0] );
magma_queue_sync( stream[d][1] );
magmablasSetKernelStream(NULL);
}
/* calling magma-gpu interface to panel-factorize the big panel */
if ( M > I ) {
magma_cgetrf2_mgpu(ngpu, M-I, N, nb, I, dAT, ldn_local, ipiv+I, dA, A(0,I), lda,
stream, &iinfo);
if ( iinfo < 0 ) {
*info = iinfo;
break;
} else if ( iinfo != 0 ) {
*info = iinfo + I * NB;
//break;
}
/* adjust pivots */
for( ii=I; ii < min(I+N,m); ii++ )
ipiv[ii] += I;
}
time_comp += timer_stop( time );
/* download the current big panel to CPU */
timer_start( time );
magmablas_cgetmatrix_transpose_mgpu(ngpu, stream, dAT, ldn_local, A(0,I), lda, dA, maxm, M, N, nb);
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_queue_sync( stream[d][0] );
magma_queue_sync( stream[d][1] );
magmablasSetKernelStream(NULL);
}
time_get += timer_stop( time );
} /* end of for */
timer_stop( time_total );
flops = FLOPS_CGETRF( m, n ) / 1e9;
timer_printf(" memory-allocation time: %e\n", time_alloc );
timer_printf(" NB=%d nb=%d\n", (int) NB, (int) nb );
timer_printf(" memcopy and transpose %e seconds\n", time_set );
timer_printf(" total time %e seconds\n", time_total );
timer_printf(" Performance %f GFlop/s, %f seconds without htod and dtoh\n", flops / (time_comp), time_comp );
timer_printf(" Performance %f GFlop/s, %f seconds with htod\n", flops / (time_comp + time_set), time_comp + time_set );
timer_printf(" Performance %f GFlop/s, %f seconds with dtoh\n", flops / (time_comp + time_get), time_comp + time_get );
timer_printf(" Performance %f GFlop/s, %f seconds without memory-allocation\n", flops / (time_total - time_alloc), time_total - time_alloc );
for( d=0; d < ngpu0; d++ ) {
magma_setdevice(d);
magma_free( dA[d] );
magma_event_destroy( event[d][0] );
magma_event_destroy( event[d][1] );
magma_queue_destroy( stream[d][0] );
magma_queue_destroy( stream[d][1] );
}
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
}
if ( *info >= 0 )
magma_cgetrf_piv(m, n, NB, A, lda, ipiv, info);
return *info;
} /* magma_cgetrf_m */
/**
Purpose
-------
CGEGQR orthogonalizes the N vectors given by a complex M-by-N matrix A:
A = Q * R.
On exit, if successful, the orthogonal vectors Q overwrite A
and R is given in work (on the CPU memory).
The routine is designed for tall-and-skinny matrices: M >> N, N <= 128.
This version uses normal equations and SVD in an iterative process that
makes the computation numerically accurate.
Arguments
---------
@param[in]
ikind INTEGER
Several versions are implemented indiceted by the ikind value:
1: This version uses normal equations and SVD in an iterative process
that makes the computation numerically accurate.
2: This version uses a standard LAPACK-based orthogonalization through
MAGMA's QR panel factorization (magma_cgeqr2x3_gpu) and magma_cungqr
3: MGS
4. Cholesky QR
@param[in]
m INTEGER
The number of rows of the matrix A. m >= n >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. 128 >= n >= 0.
@param[in,out]
dA COMPLEX array on the GPU, dimension (ldda,n)
On entry, the m-by-n matrix A.
On exit, the m-by-n matrix Q with orthogonal columns.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,m).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param
dwork (GPU workspace) COMPLEX array, dimension:
n^2 for ikind = 1
3 n^2 + min(m, n) for ikind = 2
0 (not used) for ikind = 3
n^2 for ikind = 4
@param[out]
work (CPU workspace) COMPLEX array, dimension 3 n^2.
On exit, work(1:n^2) holds the rectangular matrix R.
Preferably, for higher performance, work should be in pinned memory.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
@ingroup magma_cgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_cgegqr_gpu( magma_int_t ikind, magma_int_t m, magma_int_t n,
magmaFloatComplex *dA, magma_int_t ldda,
magmaFloatComplex *dwork, magmaFloatComplex *work,
magma_int_t *info )
{
#define work(i_,j_) (work + (i_) + (j_)*n)
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
magma_int_t i = 0, j, k, n2 = n*n;
magma_int_t ione = 1;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
float cn = 200., mins, maxs;
/* check arguments */
*info = 0;
if (ikind < 1 || ikind > 4) {
*info = -1;
} else if (m < 0 || m < n) {
*info = -2;
} else if (n < 0 || n > 128) {
*info = -3;
} else if (ldda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (ikind == 1) {
// === Iterative, based on SVD ============================================================
magmaFloatComplex *U, *VT, *vt, *R, *G, *hwork, *tau;
float *S;
R = work; // Size n * n
G = R + n*n; // Size n * n
VT = G + n*n; // Size n * n
magma_cmalloc_cpu( &hwork, 32 + 2*n*n + 2*n);
if ( hwork == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_int_t lwork=n*n+32; // First part f hwork; used as workspace in svd
U = hwork + n*n + 32; // Size n*n
S = (float *)(U+n*n); // Size n
tau = U + n*n + n; // Size n
#if defined(PRECISION_c) || defined(PRECISION_z)
float *rwork;
magma_smalloc_cpu( &rwork, 5*n);
if ( rwork == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
#endif
do {
i++;
magma_cgemm(MagmaConjTrans, MagmaNoTrans, n, n, m, c_one, dA, ldda, dA, ldda, c_zero, dwork, n );
magma_cgetmatrix(n, n, dwork, n, G, n);
#if defined(PRECISION_s) || defined(PRECISION_d)
lapackf77_cgesvd("n", "a", &n, &n, G, &n, S, U, &n, VT, &n,
hwork, &lwork, info);
#else
lapackf77_cgesvd("n", "a", &n, &n, G, &n, S, U, &n, VT, &n,
hwork, &lwork, rwork, info);
#endif
mins = 100.f, maxs = 0.f;
for (k=0; k < n; k++) {
S[k] = magma_ssqrt( S[k] );
if (S[k] < mins) mins = S[k];
if (S[k] > maxs) maxs = S[k];
}
for (k=0; k < n; k++) {
vt = VT + k*n;
for (j=0; j < n; j++)
vt[j] *= S[j];
}
lapackf77_cgeqrf(&n, &n, VT, &n, tau, hwork, &lwork, info);
if (i == 1)
blasf77_ccopy(&n2, VT, &ione, R, &ione);
else
blasf77_ctrmm("l", "u", "n", "n", &n, &n, &c_one, VT, &n, R, &n);
magma_csetmatrix(n, n, VT, n, dwork, n);
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit, m, n, c_one, dwork, n, dA, ldda);
if (mins > 0.00001f)
cn = maxs/mins;
//fprintf(stderr, "Iteration %d, cond num = %f \n", i, cn);
} while (cn > 10.f);
magma_free_cpu( hwork );
#if defined(PRECISION_c) || defined(PRECISION_z)
magma_free_cpu( rwork );
#endif
// ================== end of ikind == 1 ===================================================
}
else if (ikind == 2) {
// ================== LAPACK based ===================================================
magma_int_t min_mn = min(m, n);
magma_int_t nb = n;
magmaFloatComplex *dtau = dwork + 2*n*n, *d_T = dwork, *ddA = dwork + n*n;
magmaFloatComplex *tau = work+n*n;
magmablas_claset( MagmaFull, n, n, c_zero, c_zero, d_T, n );
magma_cgeqr2x3_gpu(&m, &n, dA, &ldda, dtau, d_T, ddA,
(float *)(dwork+min_mn+2*n*n), info);
magma_cgetmatrix( min_mn, 1, dtau, min_mn, tau, min_mn);
magma_cgetmatrix( n, n, ddA, n, work, n);
magma_cungqr_gpu( m, n, n, dA, ldda, tau, d_T, nb, info );
// ================== end of ikind == 2 ===================================================
}
else if (ikind == 3) {
// ================== MGS ===================================================
for(magma_int_t j = 0; j<n; j++){
for(magma_int_t i = 0; i<j; i++){
*work(i, j) = magma_cdotc(m, dA(0,i), 1, dA(0,j), 1);
magma_caxpy(m, -(*work(i,j)), dA(0,i), 1, dA(0,j), 1);
}
for(magma_int_t i = j; i<n; i++)
*work(i, j) = MAGMA_C_ZERO;
//*work(j,j) = MAGMA_C_MAKE( magma_scnrm2(m, dA(0,j), 1), 0. );
*work(j,j) = magma_cdotc(m, dA(0,j), 1, dA(0,j), 1);
*work(j,j) = MAGMA_C_MAKE( sqrt(MAGMA_C_REAL( *work(j,j) )), 0.);
magma_cscal(m, 1./ *work(j,j), dA(0,j), 1);
}
// ================== end of ikind == 3 ===================================================
}
else if (ikind == 4) {
// ================== Cholesky QR ===================================================
magma_cgemm(MagmaConjTrans, MagmaNoTrans, n, n, m, c_one, dA, ldda, dA, ldda, c_zero, dwork, n );
magma_cgetmatrix(n, n, dwork, n, work, n);
lapackf77_cpotrf("u", &n, work, &n, info);
magma_csetmatrix(n, n, work, n, dwork, n);
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit, m, n, c_one, dwork, n, dA, ldda);
// ================== end of ikind == 4 ===================================================
}
return *info;
} /* magma_cgegqr_gpu */
extern "C" magma_int_t
magma_chetrd_he2hb_mgpu( char uplo, magma_int_t n, magma_int_t nb,
magmaFloatComplex *a, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *work, magma_int_t lwork,
magmaFloatComplex *dAmgpu[], magma_int_t ldda,
magmaFloatComplex *dTmgpu[], magma_int_t lddt,
magma_int_t ngpu, magma_int_t distblk,
magma_queue_t streams[][20], magma_int_t nstream,
magma_int_t threads, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CHETRD_HE2HB reduces a complex Hermitian matrix A to real symmetric
band-diagonal form T by an orthogonal similarity transformation:
Q**H * A * Q = T.
This version stores the triangular matrices T used in the accumulated
Householder transformations (I - V T V').
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
A (input/output) COMPLEX array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = 'U', the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if UPLO = 'U', the Upper band-diagonal of A is
overwritten by the corresponding elements of the
band-diagonal matrix T, and the elements above the band
diagonal, with the array TAU, represent the orthogonal
matrix Q as a product of elementary reflectors; if UPLO
= 'L', the the Lower band-diagonal of A is overwritten by
the corresponding elements of the band-diagonal
matrix T, and the elements below the band-diagonal, with
the array TAU, represent the orthogonal matrix Q as a product
of elementary reflectors. See Further Details.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
TAU (output) COMPLEX array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details).
WORK (workspace/output) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The dimension of the array WORK. LWORK >= 1.
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
dT (output) COMPLEX array on the GPU, dimension N*NB,
where NB is the optimal blocksize.
On exit dT holds the upper triangular matrices T from the
accumulated Householder transformations (I - V T V') used
in the factorization. The nb x nb matrices T are ordered
consecutively in memory one after another.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
Further Details
===============
If UPLO = 'U', the matrix Q is represented as a product of elementary
reflectors
Q = H(n-1) . . . H(2) H(1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(i+1:n) = 0 and v(i) = 1; v(1:i-1) is stored on exit in
A(1:i-1,i+1), and tau in TAU(i).
If UPLO = 'L', the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(n-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0 and v(i+1) = 1; v(i+2:n) is stored on exit in A(i+2:n,i),
and tau in TAU(i).
The contents of A on exit are illustrated by the following examples
with n = 5:
if UPLO = 'U': if UPLO = 'L':
( d e v2 v3 v4 ) ( d )
( d e v3 v4 ) ( e d )
( d e v4 ) ( v1 e d )
( d e ) ( v1 v2 e d )
( d ) ( v1 v2 v3 e d )
where d and e denote diagonal and off-diagonal elements of T, and vi
denotes an element of the vector defining H(i).
===================================================================== */
#define a_ref(a_1,a_2) ( a + ((a_2)-1)*( lda) + (a_1)-1)
#define da_ref(a_1,a_2) (da + ((a_2)-1)*(ldda) + (a_1)-1)
#define tau_ref(a_1) (tau + (a_1)-1)
#define t_ref(a_1) (dT + ((a_1)-1)*(lddt))
#define Atest(a_1,a_2) ( Atest + ((a_2)-1)*( lda) + (a_1)-1)
#define dttest(a_0, a_1, a_2) (dTmgpu[a_0] + ((a_2)-1)*(lddt))
#define datest(a_0, a_1, a_2) (dAmgpu[a_0] + ((a_2)-1)*(ldda) + (a_1)-1)
char uplo_[2] = {uplo, 0};
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex c_neg_half = MAGMA_C_NEG_HALF;
magmaFloatComplex c_one = MAGMA_C_ONE ;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
float d_one = MAGMA_D_ONE;
magma_int_t pm, pn, indi, indj, pk;
magma_int_t pm_old=0, pn_old=0, indi_old=0, indj_old=0, flipV=-1;
magma_int_t iblock, idev, di;
int i;
int lwkopt;
int lquery;
assert (nstream>=3);
assert (nstream>=(ngpu+1));
*info = 0;
int upper = lapackf77_lsame(uplo_, "U");
lquery = lwork == -1;
if (! upper && ! lapackf77_lsame(uplo_, "L")) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
} else if (lwork < 1 && ! lquery) {
*info = -9;
}
/* Determine the block size. */
lwkopt = n * nb;
if (*info == 0) {
MAGMA_C_SET2REAL( work[0], lwkopt );
}
if (*info != 0)
return *info;
else if (lquery)
return *info;
/* Quick return if possible */
if (n == 0) {
work[0] = c_one;
return *info;
}
magma_int_t mklth = min(threads,16);
magma_setlapack_numthreads(mklth);
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2];
magma_int_t nbcmplx=0;
magma_buildconnection_mgpu(gnode, &nbcmplx, ngpu);
#ifdef ENABLE_DEBUG
printf(" Initializing communication pattern.... GPU-ncmplx %d\n\n" , nbcmplx);
#endif
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_t cstream;
magmablasGetKernelStream(&cstream);
magmaFloatComplex *dspace[MagmaMaxGPUs];
magmaFloatComplex *dwork[MagmaMaxGPUs], *dworkbis[MagmaMaxGPUs];
magmaFloatComplex *dvall[MagmaMaxGPUs], *dv[MagmaMaxGPUs], *dw[MagmaMaxGPUs];
magmaFloatComplex *workngpu[MagmaMaxGPUs+1];
magma_event_t redevents[MagmaMaxGPUs][MagmaMaxGPUs*MagmaMaxGPUs+10];
magma_int_t nbevents = MagmaMaxGPUs*MagmaMaxGPUs;
magma_int_t lddv = ldda;
magma_int_t lddw = lddv;
magma_int_t dwrk2siz = ldda*nb*(ngpu+1);
magma_int_t worksiz = n*nb;
magma_int_t devworksiz = 2*nb*lddv + nb*lddw + nb*ldda + dwrk2siz; // 2*dv(dv0+dv1) + dw + dwork +dworkbis
// local allocation and stream creation
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_cmalloc( &dspace[dev], devworksiz );
magma_cmalloc_pinned ( &workngpu[dev], worksiz);
dvall[dev] = dspace[dev];
dw[dev] = dvall[dev] + 2*nb*lddv;
dwork[dev] = dw[dev] + nb*lddw;
dworkbis[dev] = dwork[dev] + nb*ldda;
magmablasSetKernelStream( streams[ dev ][ 0 ] );
for( magma_int_t i = 0; i < nbevents; ++i ) {
cudaEventCreateWithFlags(&redevents[dev][i],cudaEventDisableTiming);
}
}
magma_cmalloc_pinned ( &workngpu[ngpu], worksiz);
magmaFloatComplex *worktest = NULL; //(magmaFloatComplex *) malloc(n*nb*sizeof(magmaFloatComplex)); // not used
// ======================
magmaFloatComplex *hT = work + lwork - nb*nb;
lwork -= nb*nb;
memset( hT, 0, nb*nb*sizeof(magmaFloatComplex));
if (upper) {
printf("CHETRD_HE2HB is not yet implemented for upper matrix storage. Exit.\n");
exit(1);
}else {
/* Reduce the lower triangle of A */
for (i = 1; i <= n-nb; i += nb)
{
indi = i+nb;
indj = i;
pm = n - i - nb + 1;
//pn = min(i+nb-1, n-nb) -i + 1;
pn = nb;
/* Get the current panel (no need for the 1st iteration) */
if (i > 1 ){
// cpanel_to_q copy the upper oof diagonal part of
// the matrix to work to be restored later. acctually
// the zero's and one's putted are not used this is only
// because we don't have a function that copy only the
// upper part of A to be restored after copying the
// lookahead panel that has been computted from GPU to CPU.
cpanel_to_q(MagmaUpper, pn-1, a_ref(i, i+1), lda, work);
// find the device who own the panel then send it to the CPU.
// below a -1 was added and then a -1 was done on di because of the fortran indexing
iblock = ((i-1) / distblk) / ngpu; // local block id
di = iblock*distblk + (i-1)%distblk; // local index in parent matrix
idev = ((i-1) / distblk) % ngpu; // device with this block
//printf("Receiving panel ofsize %d %d from idev %d A(%d,%d) \n",(pm+pn), pn,idev,i-1,di);
magma_setdevice( idev );
//magma_device_sync();
magma_cgetmatrix_async( (pm+pn), pn,
datest(idev, i, di+1), ldda,
a_ref ( i, i), lda, streams[ idev ][ nstream-1 ] );
/*
magma_device_sync();
cudaMemcpy2DAsync(a_ref(i,i), lda*sizeof(magmaFloatComplex),
datest(idev,i,di+1), ldda*sizeof(magmaFloatComplex),
(pm+pn)*sizeof(magmaFloatComplex), pn,
cudaMemcpyDeviceToHost, streams[ idev ][ nstream-1 ]);
*/
//magma_setdevice( 0 );
//printf("updating cher2k on A(%d,%d) of size %d %d \n",indi_old+pn_old-1,indi_old+pn_old-1,pm_old-pn_old,pn_old);
// compute CHER2K_MGPU
magmablas_cher2k_mgpu2(
MagmaLower, MagmaNoTrans, pm_old-pn_old, pn_old,
c_neg_one, dv, pm_old, pn_old,
dw, pm_old, pn_old,
d_one, dAmgpu, ldda, indi_old+pn_old-1,
ngpu, distblk, streams, 2 );
//magma_setdevice( 0 );
magma_setdevice( idev );
magma_queue_sync( streams[idev][ nstream-1 ] );
//magma_setdevice( 0 );
cq_to_panel(MagmaUpper, pn-1, a_ref(i, i+1), lda, work);
}
/* ==========================================================
QR factorization on a panel starting nb off of the diagonal.
Prepare the V and T matrices.
========================================================== */
lapackf77_cgeqrf(&pm, &pn, a_ref(indi, indj), &lda,
tau_ref(i), work, &lwork, info);
/* Form the matrix T */
pk=min(pm,pn);
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&pm, &pk, a_ref(indi, indj), &lda,
tau_ref(i), hT, &nb);
/* Prepare V - put 0s in the upper triangular part of the panel
(and 1s on the diagonal), temporaly storing the original in work */
cpanel_to_q(MagmaUpper, pk, a_ref(indi, indj), lda, work);
/* Send V and T from the CPU to the GPU */
// To be able to overlap the GET with the CHER2K
// it should be done on last stream.
// TO Avoid a BUG that is overwriting the old_V
// used atthis moment by cher2k with the new_V
// send it now, we decide to have a flipflop
// vector of Vs. if step%2=0 use V[0] else use V[nb*n]
flipV = ((i-1)/nb)%2;
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
dv[dev] = dvall[dev] + flipV*nb*lddv;
}
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
// send V
magma_csetmatrix_async( pm, pk,
a_ref(indi, indj), lda,
dv[dev], pm, streams[dev][nstream-1] );
// Send the triangular factor T to the GPU
magma_csetmatrix_async( pk, pk,
hT, nb,
dttest(dev, 1, i), lddt, streams[dev][nstream-1] );
}
/* ==========================================================
Compute W:
1. X = A (V T)
2. W = X - 0.5* V * (T' * (V' * X))
========================================================== */
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
// dwork = V T
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ nstream-1 ] );
magma_queue_sync( streams[dev][nstream-1] );
magma_cgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_one, dv[dev], pm,
dttest(dev, 1, i), lddt,
c_zero, dwork[dev], pm);
}
// ===============================================
// SYNC TO BE SURE THAT BOTH V AND T WERE
// RECEIVED AND VT IS COMPUTED and SYR2K is done
// ===============================================
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
for( magma_int_t s = 0; s < nstream; ++s )
magma_queue_sync( streams[dev][s] );
}
// compute CHEMM_MGPU
// The broadcast of the result done inside this function
// should be done in stream [0] because i am assuming this
// for the GEMMs below otherwise I have to SYNC over the
// Broadcasting stream.
if(ngpu==1){
magmablasSetKernelStream( streams[ 0 ][ 0 ] );
magma_chemm(MagmaLeft, uplo, pm, pk,
c_one, dAmgpu[0]+(indi-1)*ldda+(indi-1), ldda,
dwork[0], pm,
c_zero, dw[0], pm);
}else{
magmablas_chemm_mgpu_com(
MagmaLeft, uplo, pm, pk,
c_one, dAmgpu, ldda, indi-1,
dwork, pm,
c_zero, dw, pm, dworkbis, dwrk2siz, worktest, pm, workngpu, worksiz,
ngpu, distblk, streams, nstream-1, redevents, nbevents, gnode, nbcmplx);
}
/* dwork = V*T already ==> dwork' = T'*V'
* compute T'*V'*X ==> dwork'*W ==>
* dwork + pm*nb = ((T' * V') * X) = dwork' * X = dwork' * W */
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
// Here we have to wait until the broadcast of CHEMM has been done.
// Note that the broadcast should be done on stream[0] so in a way
// we can continue here on the same stream and avoid a sync
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
// magma_queue_sync( streams[dev][0] );
magma_cgemm(MagmaConjTrans, MagmaNoTrans, pk, pk, pm,
c_one, dwork[dev], pm,
dw[dev], pm,
c_zero, dworkbis[dev], nb);
/* W = X - 0.5 * V * T'*V'*X
* = X - 0.5 * V * (dwork + pm*nb) = W - 0.5 * V * (dwork + pm*nb) */
magma_cgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_neg_half, dv[dev], pm,
dworkbis[dev], nb,
c_one, dw[dev], pm);
}
/* restore the panel it is put here to overlap with the previous GEMM*/
cq_to_panel(MagmaUpper, pk, a_ref(indi, indj), lda, work);
// ===============================================
// SYNC TO BE SURE THAT BOTH V AND W ARE DONE
// ===============================================
// Synchronise to be sure that W has been computed
// because next CHER2K use streaming and may happen
// that lunch a gemm on stream 2 while stream 0
// which compute those 2 GEMM above has not been
// computed and also used for the same reason in
// the panel update below and also for the last HER2K
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_queue_sync( streams[dev][0] );
}
/* ==========================================================
Update the unreduced submatrix A(i+ib:n,i+ib:n), using
an update of the form: A := A - V*W' - W*V'
========================================================== */
if (i + nb <= n-nb){
/* There would be next iteration;
do lookahead - update the next panel */
// below a -1 was added and then a -1 was done on di because of the fortran indexing
iblock = ((indi-1) / distblk) / ngpu; // local block id
di = iblock*distblk + (indi-1)%distblk; // local index in parent matrix
idev = ((indi-1) / distblk) % ngpu; // device with this block
magma_setdevice( idev );
magmablasSetKernelStream( streams[ idev ][ nstream-1 ] );
//magma_queue_sync( streams[idev][0] ); removed because the sync has been done in the loop above
magma_cgemm(MagmaNoTrans, MagmaConjTrans, pm, pn, pn, c_neg_one,
dv[idev], pm,
dw[idev] , pm, c_one,
datest(idev, indi, di+1), ldda);
magma_cgemm(MagmaNoTrans, MagmaConjTrans, pm, pn, pn, c_neg_one,
dw[idev] , pm,
dv[idev], pm, c_one,
datest(idev, indi, di+1), ldda);
//printf("updating next panel distblk %d idev %d on A(%d,%d) of size %d %d %d \n",distblk,idev,indi-1,di,pm,pn,pn);
}
else {
/* no look-ahead as this is last iteration */
// below a -1 was added and then a -1 was done on di because of the fortran indexing
iblock = ((indi-1) / distblk) / ngpu; // local block id
di = iblock*distblk + (indi-1)%distblk; // local index in parent matrix
idev = ((indi-1) / distblk) % ngpu; // device with this block
magma_setdevice( idev );
magmablasSetKernelStream( streams[ idev ][ 0 ] );
//printf("LAST CHER2K idev %d on A(%d,%d) of size %d \n",idev, indi-1,di,pk);
magma_cher2k(MagmaLower, MagmaNoTrans, pk, pk, c_neg_one,
dv[idev], pm,
dw[idev] , pm, d_one,
datest(idev, indi, di+1), ldda);
/* Send the last block to the CPU */
cpanel_to_q(MagmaUpper, pk-1, a_ref(n-pk+1, n-pk+2), lda, work);
magma_cgetmatrix( pk, pk,
datest(idev, indi, di+1), ldda,
a_ref(n-pk+1, n-pk+1), lda );
cq_to_panel(MagmaUpper, pk-1, a_ref(n-pk+1, n-pk+2), lda, work);
}
indi_old = indi;
indj_old = indj;
pm_old = pm;
pn_old = pn;
} // end loop for(i)
}// end of LOWER
//magma_setdevice( 0 );
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_free( dspace[dev]);
magma_free_pinned(workngpu[dev]);
for( magma_int_t e = 0; e < nbevents; ++e ) {
cudaEventDestroy(redevents[dev][e]);
}
}
magma_free_pinned(workngpu[ngpu]);
free(worktest);
magma_setdevice( cdev );
magmablasSetKernelStream( cstream );
MAGMA_C_SET2REAL( work[0], lwkopt );
magma_setlapack_numthreads(1);
return *info;
} /* chetrd_he2hb_ */
extern "C" void
magma_cherk_mgpu(
magma_int_t ngpu,
magma_uplo_t uplo, magma_trans_t trans, magma_int_t nb, magma_int_t n, magma_int_t k,
float alpha,
magmaFloatComplex_ptr dB[], magma_int_t lddb, magma_int_t b_offset,
float beta,
magmaFloatComplex_ptr dC[], magma_int_t lddc, magma_int_t c_offset,
magma_int_t nqueue, magma_queue_t queues[][10])
{
#define dB(id, i, j) (dB[(id)]+(j)*lddb + (i)+b_offset)
#define dC(id, i, j) (dC[(id)]+(j)*lddc + (i))
#define STREAM_ID(i) (nqueue > 1 ? 1+((i)/nb)%(nqueue-1) : 0)
magma_int_t i, id, ib, ii, kk, n1;
magmaFloatComplex z_alpha = MAGMA_C_MAKE(alpha,0.0);
magmaFloatComplex z_beta = MAGMA_C_MAKE(beta, 0.0);
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
/* diagonal update */
for( i=0; i < n; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = STREAM_ID( i+c_offset );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
/* cher2k on diagonal block */
magma_setdevice(id);
magmablasSetKernelStream( queues[id][kk] );
trace_gpu_start( id, kk, "syr2k", "syr2k" );
magma_cherk(uplo, trans, ib, k,
alpha, dB(id, i, 0 ), lddb,
beta, dC(id, i+c_offset, ii), lddc);
trace_gpu_end( id, kk );
}
/* off-diagonal update */
if (uplo == MagmaUpper) {
for( i=nb; i < n; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = STREAM_ID( i+c_offset );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
magma_setdevice(id);
magmablasSetKernelStream( queues[id][kk] );
magma_cgemm(MagmaNoTrans, MagmaConjTrans, i, ib, k,
z_alpha, dB(id, 0, 0 ), lddb,
dB(id, i, 0 ), lddb,
z_beta, dC(id, 0, ii), lddc);
}
}
else {
for( i=0; i < n-nb; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = STREAM_ID( i+c_offset );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
n1 = n-i-ib;
/* cgemm on off-diagonal blocks */
magma_setdevice(id);
magmablasSetKernelStream( queues[id][kk] );
trace_gpu_start( id, kk, "gemm_up", "gemm_up" );
magma_cgemm(MagmaNoTrans, MagmaConjTrans, n1, ib, k,
z_alpha, dB(id, i+ib, 0 ), lddb,
dB(id, i, 0 ), lddb,
z_beta, dC(id, i+c_offset+ib, ii), lddc);
trace_gpu_end( id, kk );
}
}
// TODO why not sync?
//for( id=0; id < ngpu; id++ ) {
// magma_setdevice(id);
// //for( kk=0; kk < nqueue; kk++ )
// // magma_queue_sync( queues[id][kk] );
//}
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
}
extern "C" magma_int_t
magma_cgessm_gpu( char storev, magma_int_t m, magma_int_t n, magma_int_t k, magma_int_t ib,
magma_int_t *ipiv,
cuFloatComplex *dL1, magma_int_t lddl1,
cuFloatComplex *dL, magma_int_t lddl,
cuFloatComplex *dA, magma_int_t ldda,
magma_int_t *info)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
CGESSM applies the factors L computed by CGETRF_INCPIV to
a complex M-by-N tile A.
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.
K (input) INTEGER
The number of columns of the matrix L. K >= 0.
IB (input) INTEGER
The inner-blocking size. IB >= 0.
IPIV (input) INTEGER array on the cpu.
The pivot indices array of size K as returned by
CGETRF_INCPIV.
dL1 (input) DOUBLE COMPLEX array, dimension(LDDL1, N)
The IB-by-K matrix in which is stored L^(-1) as returned by GETRF_INCPIV
LDDL1 (input) INTEGER
The leading dimension of the array L1. LDDL1 >= max(1,2*IB).
dL (input) DOUBLE COMPLEX array, dimension(LDDL, N)
The M-by-K lower triangular tile on the gpu.
LDDL (input) INTEGER
The leading dimension of the array L. LDDL >= max(1,M).
dA (input/output) DOUBLE COMPLEX array, dimension (LDDA, N)
On entry, the M-by-N tile A on the gpu.
On exit, updated by the application of L on the gpu.
===================================================================== */
#define AT(i,j) (dAT + (i)*ldda + (j) )
#define L(i,j) (dL + (i) + (j)*lddl )
#define dL1(j) (dL1 + (j)*lddl1)
cuFloatComplex c_one = MAGMA_C_ONE;
cuFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
int i, s, sb;
cuFloatComplex *dAT;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
if ( (storev == 'C') || (storev == 'c') ) {
magmablas_cgetmo_in( dA, dAT, ldda, m, n );
} else {
dAT = dA;
}
s = k / ib;
for(i = 0; i < k; i += ib) {
sb = min(ib, k-i);
magmablas_claswp( n, dAT, ldda, i+1, i+sb, ipiv, 1 );
#ifndef WITHOUTTRTRI
magma_ctrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n, sb,
c_one, dL1(i), lddl1,
AT(i, 0), ldda);
#else
magma_ctrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n, sb,
c_one, L( i, i), lddl,
AT(i, 0), ldda);
#endif
if ( (i+sb) < m) {
magma_cgemm( MagmaNoTrans, MagmaTrans,
n, m-(i+sb), sb,
c_neg_one, AT(i, 0), ldda,
L( i+sb, i), lddl,
c_one, AT(i+sb, 0), ldda );
}
}
if ( (storev == 'C') || (storev == 'c') ) {
magmablas_cgetmo_in( dA, dAT, ldda, m, n );
}
return *info;
/* End of MAGMA_CGETRF_GPU */
}
/**
@deprecated
Purpose
-------
CLAQPS 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 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 array, dimension (KB)
The scalar factors of the elementary reflectors.
@param[in,out]
vn1 REAL array, dimension (N)
The vector with the partial column norms.
@param[in,out]
vn2 REAL array, dimension (N)
The vector with the exact column norms.
@param[in,out]
dauxv COMPLEX array, dimension (NB), on the GPU
Auxiliary vector.
@param[in,out]
dF COMPLEX 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_cgeqp3_aux
********************************************************************/
extern "C" magma_int_t
magma_claqps_gpu(
magma_int_t m, magma_int_t n, magma_int_t offset,
magma_int_t nb, magma_int_t *kb,
magmaFloatComplex_ptr dA, magma_int_t ldda,
magma_int_t *jpvt, magmaFloatComplex *tau,
float *vn1, float *vn2,
magmaFloatComplex_ptr dauxv,
magmaFloatComplex_ptr dF, magma_int_t lddf)
{
#define dA(i, j) (dA + (i) + (j)*(ldda))
#define dF(i, j) (dF + (i) + (j)*(lddf))
magmaFloatComplex c_zero = MAGMA_C_MAKE( 0.,0.);
magmaFloatComplex c_one = MAGMA_C_MAKE( 1.,0.);
magmaFloatComplex c_neg_one = MAGMA_C_MAKE(-1.,0.);
magma_int_t ione = 1;
magma_int_t i__1, i__2;
magmaFloatComplex z__1;
magma_int_t k, rk;
magmaFloatComplex_ptr dAks;
magmaFloatComplex tauk = MAGMA_C_ZERO;
magma_int_t pvt;
float tol3z;
magma_int_t itemp;
float lsticc;
magmaFloat_ptr dlsticcs;
magma_smalloc( &dlsticcs, 1+256*(n+255)/256 );
tol3z = magma_ssqrt( lapackf77_slamch("Epsilon"));
lsticc = 0;
k = 0;
magma_cmalloc( &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 isamax; pvt, k are 0-based.
pvt = k + magma_isamax( 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_cswap( m, dA(0, pvt), ione, dA(0, k), ione, queue );
magmablas_cswap( i__1, dF(pvt, 0), lddf, dF(k, 0), lddf, queue );
itemp = jpvt[pvt];
jpvt[pvt] = jpvt[k];
jpvt[k] = itemp;
magma_sswap( 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_cgemv( 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_cgemv( 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_clarfg_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_csetvector( 1, &c_one, 1, dA(rk, k), 1, queue );
else magma_ccopymatrix( 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_cgetvector( 1, &tau[k], 1, &tauk, 1, queue );
if (k < n-1) {
i__1 = m - rk;
i__2 = n - k - 1;
/* Multiply on GPU */
magma_cgemv( 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_C_NEGATE( tauk );
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_cgemv( 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_cgemv( 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_cgemv( 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_cgemv( 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_cgemm( 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_cgemm( 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_scnrm2_row_check_adjust( n-k-1, tol3z, &vn1[k+1], &vn2[k+1],
dA(rk,k+1), ldda, dlsticcs, queue );
//magma_device_sync();
magma_sgetvector( 1, &dlsticcs[0], 1, &lsticc, 1, queue );
}
++k;
}
magma_ccopymatrix( 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_cgemm( 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_scnrm2_check( m-rk-1, n-*kb, dA(rk+1,*kb), ldda,
&vn1[*kb], dlsticcs, queue );
magma_scopymatrix( n-*kb, 1, &vn1[*kb], *kb, &vn2[*kb], *kb, queue );
}
magma_free( dAks );
magma_free( dlsticcs );
magma_queue_destroy( queue );
return MAGMA_SUCCESS;
} /* magma_claqps */
/**
Purpose
-------
CGESSM applies the factors L computed by CGETRF_INCPIV to
a complex M-by-N tile A.
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]
k INTEGER
The number of columns of the matrix L. K >= 0.
@param[in]
ib INTEGER
The inner-blocking size. IB >= 0.
@param[in]
ipiv INTEGER array on the cpu.
The pivot indices array of size K as returned by
CGETRF_INCPIV.
@param[in]
dL1 COMPLEX array, dimension(LDDL1, N)
The IB-by-K matrix in which is stored L^(-1) as returned by GETRF_INCPIV
@param[in]
lddl1 INTEGER
The leading dimension of the array L1. LDDL1 >= max(1,2*IB).
@param[in]
dL COMPLEX array, dimension(LDDL, N)
The M-by-K lower triangular tile on the gpu.
@param[in]
lddl INTEGER
The leading dimension of the array L. LDDL >= max(1,M).
@param[in,out]
dA COMPLEX array, dimension (LDDA, N)
On entry, the M-by-N tile A on the gpu.
On exit, updated by the application of L on the gpu.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
@ingroup magma_cgesv_tile
********************************************************************/
extern "C" magma_int_t
magma_cgessm_gpu(
magma_order_t order, magma_int_t m, magma_int_t n, magma_int_t k, magma_int_t ib,
magma_int_t *ipiv,
magmaFloatComplex_ptr dL1, magma_int_t lddl1,
magmaFloatComplex_ptr dL, magma_int_t lddl,
magmaFloatComplex_ptr dA, magma_int_t ldda,
magma_int_t *info)
{
#define AT(i,j) (dAT + (i)*ldda + (j) )
#define L(i,j) (dL + (i) + (j)*lddl )
#define dL1(j) (dL1 + (j)*lddl1)
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
int i, sb;
magmaFloatComplex_ptr dAT;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
if ( order == MagmaColMajor ) {
magmablas_cgetmo_in( dA, dAT, ldda, m, n );
} else {
dAT = dA;
}
for (i = 0; i < k; i += ib) {
sb = min(ib, k-i);
magmablas_claswp( n, dAT, ldda, i+1, i+sb, ipiv, 1 );
#ifndef WITHOUTTRTRI
magma_ctrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n, sb,
c_one, dL1(i), lddl1,
AT(i, 0), ldda);
#else
magma_ctrsm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n, sb,
c_one, L( i, i), lddl,
AT(i, 0), ldda);
#endif
if ( (i+sb) < m) {
magma_cgemm( MagmaNoTrans, MagmaTrans,
n, m-(i+sb), sb,
c_neg_one, AT(i, 0), ldda,
L( i+sb, i), lddl,
c_one, AT(i+sb, 0), ldda );
}
}
if ( order == MagmaColMajor ) {
magmablas_cgetmo_in( dA, dAT, ldda, m, n );
}
return *info;
} /* magma_cgessm_gpu */
// ----------------------------------------------------------------------
extern "C" void
magma_cher2k_mgpu(
magma_int_t ngpu,
magma_uplo_t uplo, magma_trans_t trans, magma_int_t nb, magma_int_t n, magma_int_t k,
magmaFloatComplex alpha,
magmaFloatComplex_ptr dB[], magma_int_t lddb, magma_int_t b_offset,
float beta,
magmaFloatComplex_ptr dC[], magma_int_t lddc, magma_int_t c_offset,
magma_int_t nqueue, magma_queue_t queues[][10])
{
#define dB(id, i, j) (dB[(id)] + (j)*lddb + (i)+b_offset)
#define dB1(id, i, j) (dB[(id)] + (j)*lddb + (i)+b_offset) + k*lddb
#define dC(id, i, j) (dC[(id)] + (j)*lddc + (i))
magma_int_t i, id, ib, ii, kk, n1;
magmaFloatComplex c_one = MAGMA_C_ONE;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
/* diagonal update */
for( i=0; i < n; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = (i/(nb*ngpu))%nqueue;
magma_setdevice( id );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
/* cher2k on diagonal block */
trace_gpu_start( id, kk, "syr2k", "syr2k" );
magma_cher2k( uplo, trans, ib, k,
alpha, dB1(id, i, 0 ), lddb,
dB(id, i, 0 ), lddb,
beta, dC(id, i+c_offset, ii), lddc, queues[id][kk] );
trace_gpu_end( id, kk );
}
/* off-diagonal update */
if (uplo == MagmaUpper) {
for( i=nb; i < n; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = (i/(nb*ngpu))%nqueue;
magma_setdevice( id );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
magma_cgemm( MagmaNoTrans, MagmaConjTrans, i, ib, k,
alpha, dB1(id, 0, 0 ), lddb,
dB(id, i, 0 ), lddb,
c_one, dC(id, 0, ii), lddc, queues[id][kk] );
}
}
else {
for( i=0; i < n-nb; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = (i/(nb*ngpu))%nqueue;
magma_setdevice( id );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
n1 = n-i-ib;
// cgemm on off-diagonal blocks
trace_gpu_start( id, kk, "gemm_up", "gemm_up" );
magma_cgemm( MagmaNoTrans, MagmaConjTrans, n1, ib, k,
alpha, dB1(id, i+ib, 0 ), lddb,
dB(id, i, 0 ), lddb,
c_one, dC(id, i+c_offset+ib, ii), lddc, queues[id][kk] );
trace_gpu_end( id, kk );
}
}
if (uplo == MagmaUpper) {
for( i=nb; i < n; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = (i/(nb*ngpu))%nqueue;
magma_setdevice( id );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
magma_cgemm( MagmaNoTrans, MagmaConjTrans, i, ib, k,
alpha, dB( id, 0, 0 ), lddb,
dB1(id, i, 0 ), lddb,
c_one, dC(id, 0, ii), lddc, queues[id][kk] );
}
} else {
for( i=0; i < n-nb; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = (i/(nb*ngpu))%nqueue;
magma_setdevice( id );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
n1 = n-i-ib;
/* cgemm on off-diagonal blocks */
trace_gpu_start( id, kk, "gemm_up", "gemm_up" );
magma_cgemm( MagmaNoTrans, MagmaConjTrans, n1, ib, k,
alpha, dB(id, i+ib, 0 ), lddb,
dB1(id, i, 0 ), lddb,
c_one, dC(id, i+c_offset+ib, ii), lddc, queues[id][kk] );
trace_gpu_end( id, kk );
}
}
for( id=0; id < ngpu; id++ ) {
magma_setdevice( id );
for( kk=0; kk < nqueue; kk++ ) {
magma_queue_sync( queues[id][kk] );
}
}
magma_setdevice( orig_dev );
}
/**
Purpose
-------
CLAUUM computes the product U * U' or L' * L, where the triangular
factor U or L is stored in the upper or lower triangular part of
the array dA.
If UPLO = MagmaUpper then the upper triangle of the result is stored,
overwriting the factor U in dA.
If UPLO = MagmaLower then the lower triangle of the result is stored,
overwriting the factor L in dA.
This is the blocked form of the algorithm, calling Level 3 BLAS.
Arguments
---------
@param[in]
uplo magma_uplo_t
Specifies whether the triangular factor stored in the array dA
is upper or lower triangular:
- = MagmaUpper: Upper triangular
- = MagmaLower: Lower triangular
@param[in]
n INTEGER
The order of the triangular factor U or L. N >= 0.
@param[in,out]
dA REAL array on the GPU, dimension (LDDA,N)
On entry, the triangular factor U or L.
On exit, if UPLO = MagmaUpper, the upper triangle of dA is
overwritten with the upper triangle of the product U * U';
if UPLO = MagmaLower, the lower triangle of dA is overwritten with
the lower triangle of the product L' * L.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -k, the k-th argument had an illegal value
@ingroup magma_cposv_aux
***************************************************************************/
extern "C" magma_int_t
magma_clauum_gpu(magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex *dA, magma_int_t ldda, magma_int_t *info)
{
#define dA(i, j) (dA + (j)*ldda + (i))
/* Local variables */
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t nb, i, ib;
float d_one = MAGMA_D_ONE;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex *work;
int upper = (uplo == MagmaUpper);
*info = 0;
if (! upper && uplo != MagmaLower)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,n))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
nb = magma_get_cpotrf_nb(n);
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, nb*nb )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
if (nb <= 1 || nb >= n) {
magma_cgetmatrix( n, n, dA, ldda, work, n );
lapackf77_clauum(uplo_, &n, work, &n, info);
magma_csetmatrix( n, n, work, n, dA, ldda );
}
else {
if (upper) {
/* Compute inverse of upper triangular matrix */
for (i=0; i < n; i += nb) {
ib = min(nb, (n-i));
/* Compute the product U * U'. */
magma_ctrmm( MagmaRight, MagmaUpper,
MagmaConjTrans, MagmaNonUnit, i, ib,
c_one, dA(i,i), ldda, dA(0, i),ldda);
magma_cgetmatrix( ib, ib,
dA(i, i), ldda,
work, ib );
lapackf77_clauum(MagmaUpperStr, &ib, work, &ib, info);
magma_csetmatrix( ib, ib,
work, ib,
dA(i, i), ldda );
if (i+ib < n) {
magma_cgemm( MagmaNoTrans, MagmaConjTrans,
i, ib, (n-i-ib), c_one, dA(0,i+ib),
ldda, dA(i, i+ib), ldda, c_one,
dA(0,i), ldda);
magma_cherk( MagmaUpper, MagmaNoTrans, ib,(n-i-ib),
d_one, dA(i, i+ib), ldda,
d_one, dA(i, i), ldda);
}
}
}
else {
/* Compute the product L' * L. */
for (i=0; i < n; i += nb) {
ib=min(nb,(n-i));
magma_ctrmm( MagmaLeft, MagmaLower,
MagmaConjTrans, MagmaNonUnit, ib,
i, c_one, dA(i,i), ldda,
dA(i, 0),ldda);
magma_cgetmatrix( ib, ib,
dA(i, i), ldda,
work, ib );
lapackf77_clauum(MagmaLowerStr, &ib, work, &ib, info);
magma_csetmatrix( ib, ib,
work, ib,
dA(i, i), ldda );
if (i+ib < n) {
magma_cgemm( MagmaConjTrans, MagmaNoTrans,
ib, i, (n-i-ib), c_one, dA( i+ib,i),
ldda, dA(i+ib, 0),ldda, c_one,
dA(i,0), ldda);
magma_cherk( MagmaLower, MagmaConjTrans, ib, (n-i-ib),
d_one, dA(i+ib, i), ldda,
d_one, dA(i, i), ldda);
}
}
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free_pinned( work );
return *info;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgemm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, dev_perf, dev_time, cpu_perf, cpu_time;
float magma_error, dev_error, Cnorm, work[1];
magma_int_t M, N, K;
magma_int_t Am, An, Bm, Bn;
magma_int_t sizeA, sizeB, sizeC;
magma_int_t lda, ldb, ldc, ldda, lddb, lddc;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magmaFloatComplex *h_A, *h_B, *h_C, *h_Cmagma, *h_Cdev;
magmaFloatComplex_ptr d_A, d_B, d_C;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex alpha = MAGMA_C_MAKE( 0.29, -0.86 );
magmaFloatComplex beta = MAGMA_C_MAKE( -0.48, 0.38 );
magma_opts opts;
opts.parse_opts( argc, argv );
float tol = opts.tolerance * lapackf77_slamch("E");
#ifdef HAVE_CUBLAS
// for CUDA, we can check MAGMA vs. CUBLAS, without running LAPACK
printf("%% If running lapack (option --lapack), MAGMA and %s error are both computed\n"
"%% relative to CPU BLAS result. Else, MAGMA error is computed relative to %s result.\n\n",
g_platform_str, g_platform_str );
printf("%% transA = %s, transB = %s\n",
lapack_trans_const(opts.transA),
lapack_trans_const(opts.transB) );
printf("%% M N K MAGMA Gflop/s (ms) %s Gflop/s (ms) CPU Gflop/s (ms) MAGMA error %s error\n",
g_platform_str, g_platform_str );
#else
// for others, we need LAPACK for check
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf("%% transA = %s, transB = %s\n",
lapack_trans_const(opts.transA),
lapack_trans_const(opts.transB) );
printf("%% M N K %s Gflop/s (ms) CPU Gflop/s (ms) %s error\n",
g_platform_str, g_platform_str );
#endif
printf("%%========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
K = opts.ksize[itest];
gflops = FLOPS_CGEMM( M, N, K ) / 1e9;
if ( opts.transA == MagmaNoTrans ) {
lda = Am = M;
An = K;
} else {
lda = Am = K;
An = M;
}
if ( opts.transB == MagmaNoTrans ) {
ldb = Bm = K;
Bn = N;
} else {
ldb = Bm = N;
Bn = K;
}
ldc = M;
ldda = magma_roundup( lda, opts.align ); // multiple of 32 by default
lddb = magma_roundup( ldb, opts.align ); // multiple of 32 by default
lddc = magma_roundup( ldc, opts.align ); // multiple of 32 by default
sizeA = lda*An;
sizeB = ldb*Bn;
sizeC = ldc*N;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, lda*An );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, ldb*Bn );
TESTING_MALLOC_CPU( h_C, magmaFloatComplex, ldc*N );
TESTING_MALLOC_CPU( h_Cmagma, magmaFloatComplex, ldc*N );
TESTING_MALLOC_CPU( h_Cdev, magmaFloatComplex, ldc*N );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*An );
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, lddb*Bn );
TESTING_MALLOC_DEV( d_C, magmaFloatComplex, lddc*N );
/* Initialize the matrices */
lapackf77_clarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_clarnv( &ione, ISEED, &sizeB, h_B );
lapackf77_clarnv( &ione, ISEED, &sizeC, h_C );
magma_csetmatrix( Am, An, h_A, lda, d_A, ldda, opts.queue );
magma_csetmatrix( Bm, Bn, h_B, ldb, d_B, lddb, opts.queue );
/* =====================================================================
Performs operation using MAGMABLAS (currently only with CUDA)
=================================================================== */
#ifdef HAVE_CUBLAS
magma_csetmatrix( M, N, h_C, ldc, d_C, lddc, opts.queue );
magma_time = magma_sync_wtime( opts.queue );
magmablas_cgemm( opts.transA, opts.transB, M, N, K,
alpha, d_A, ldda,
d_B, lddb,
beta, d_C, lddc,
opts.queue );
magma_time = magma_sync_wtime( opts.queue ) - magma_time;
magma_perf = gflops / magma_time;
magma_cgetmatrix( M, N, d_C, lddc, h_Cmagma, ldc, opts.queue );
#endif
/* =====================================================================
Performs operation using CUBLAS / clBLAS / Xeon Phi MKL
=================================================================== */
magma_csetmatrix( M, N, h_C, ldc, d_C, lddc, opts.queue );
dev_time = magma_sync_wtime( opts.queue );
#ifdef HAVE_CUBLAS
// opts.handle also uses opts.queue
cublasCgemm( opts.handle, cublas_trans_const(opts.transA), cublas_trans_const(opts.transB), M, N, K,
&alpha, d_A, ldda,
d_B, lddb,
&beta, d_C, lddc );
#else
magma_cgemm( opts.transA, opts.transB, M, N, K,
alpha, d_A, 0, ldda,
d_B, 0, lddb,
beta, d_C, 0, lddc, opts.queue );
#endif
dev_time = magma_sync_wtime( opts.queue ) - dev_time;
dev_perf = gflops / dev_time;
magma_cgetmatrix( M, N, d_C, lddc, h_Cdev, ldc, opts.queue );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_cgemm( lapack_trans_const(opts.transA), lapack_trans_const(opts.transB), &M, &N, &K,
&alpha, h_A, &lda,
h_B, &ldb,
&beta, h_C, &ldc );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
if ( opts.lapack ) {
// compute relative error for both magma & dev, relative to lapack,
// |C_magma - C_lapack| / |C_lapack|
Cnorm = lapackf77_clange( "F", &M, &N, h_C, &ldc, work );
blasf77_caxpy( &sizeC, &c_neg_one, h_C, &ione, h_Cdev, &ione );
dev_error = lapackf77_clange( "F", &M, &N, h_Cdev, &ldc, work ) / Cnorm;
#ifdef HAVE_CUBLAS
blasf77_caxpy( &sizeC, &c_neg_one, h_C, &ione, h_Cmagma, &ione );
magma_error = lapackf77_clange( "F", &M, &N, h_Cmagma, &ldc, work ) / Cnorm;
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) M, (int) N, (int) K,
magma_perf, 1000.*magma_time,
dev_perf, 1000.*dev_time,
cpu_perf, 1000.*cpu_time,
magma_error, dev_error,
(magma_error < tol && dev_error < tol ? "ok" : "failed"));
status += ! (magma_error < tol && dev_error < tol);
#else
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) M, (int) N, (int) K,
dev_perf, 1000.*dev_time,
cpu_perf, 1000.*cpu_time,
dev_error,
(dev_error < tol ? "ok" : "failed"));
status += ! (dev_error < tol);
#endif
}
else {
#ifdef HAVE_CUBLAS
// compute relative error for magma, relative to dev (currently only with CUDA)
Cnorm = lapackf77_clange( "F", &M, &N, h_Cdev, &ldc, work );
blasf77_caxpy( &sizeC, &c_neg_one, h_Cdev, &ione, h_Cmagma, &ione );
magma_error = lapackf77_clange( "F", &M, &N, h_Cmagma, &ldc, work ) / Cnorm;
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) --- ( --- ) %8.2e --- %s\n",
(int) M, (int) N, (int) K,
magma_perf, 1000.*magma_time,
dev_perf, 1000.*dev_time,
magma_error,
(magma_error < tol ? "ok" : "failed"));
status += ! (magma_error < tol);
#else
printf("%5d %5d %5d %7.2f (%7.2f) --- ( --- ) ---\n",
(int) M, (int) N, (int) K,
dev_perf, 1000.*dev_time );
#endif
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_C );
TESTING_FREE_CPU( h_Cmagma );
TESTING_FREE_CPU( h_Cdev );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
TESTING_FREE_DEV( d_C );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/***************************************************************************//**
Purpose
-------
CGEQRS solves the least squares problem
min || A*X - C ||
using the QR factorization A = Q*R computed by CGEQRF_GPU.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. M >= N >= 0.
@param[in]
nrhs INTEGER
The number of columns of the matrix C. NRHS >= 0.
@param[in]
dA COMPLEX array on the GPU, dimension (LDDA,N)
The i-th column must contain the vector which defines the
elementary reflector H(i), for i = 1,2,...,n, as returned by
CGEQRF_GPU in the first n columns of its array argument A.
@param[in]
ldda INTEGER
The leading dimension of the array A, LDDA >= M.
@param[in]
tau COMPLEX array, dimension (N)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by MAGMA_CGEQRF_GPU.
@param[in,out]
dB COMPLEX array on the GPU, dimension (LDDB,NRHS)
On entry, the M-by-NRHS matrix C.
On exit, the N-by-NRHS solution matrix X.
@param[in,out]
dT COMPLEX array that is the output (the 6th argument)
of magma_cgeqrf_gpu of size
2*MIN(M, N)*NB + ceil(N/32)*32 )* MAX(NB, NRHS).
The array starts with a block of size MIN(M,N)*NB that stores
the triangular T matrices used in the QR factorization,
followed by MIN(M,N)*NB block storing the diagonal block
inverses for the R matrix, followed by work space of size
(ceil(N/32)*32)* MAX(NB, NRHS).
@param[in]
lddb INTEGER
The leading dimension of the array dB. LDDB >= M.
@param[out]
hwork (workspace) COMPLEX array, dimension (LWORK)
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK,
LWORK >= (M - N + NB)*(NRHS + NB) + NRHS*NB,
where NB is the blocksize given by magma_get_cgeqrf_nb( M, N ).
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the HWORK array, returns
this value as the first entry of the WORK array.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_geqrs
*******************************************************************************/
extern "C" magma_int_t
magma_cgeqrs_gpu(
magma_int_t m, magma_int_t n, magma_int_t nrhs,
magmaFloatComplex_const_ptr dA, magma_int_t ldda,
magmaFloatComplex const *tau,
magmaFloatComplex_ptr dT,
magmaFloatComplex_ptr dB, magma_int_t lddb,
magmaFloatComplex *hwork, magma_int_t lwork,
magma_int_t *info)
{
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
#define dT(i_) (dT + (lddwork + (i_))*nb)
/* Constants */
const magmaFloatComplex c_zero = MAGMA_C_ZERO;
const magmaFloatComplex c_one = MAGMA_C_ONE;
const magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
const magma_int_t ione = 1;
/* Local variables */
magmaFloatComplex_ptr dwork;
magma_int_t i, min_mn, lddwork, rows, ib;
magma_int_t nb = magma_get_cgeqrf_nb( m, n );
magma_int_t lwkopt = (m - n + nb)*(nrhs + nb) + nrhs*nb;
bool lquery = (lwork == -1);
hwork[0] = magma_cmake_lwork( lwkopt );
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0 || m < n)
*info = -2;
else if (nrhs < 0)
*info = -3;
else if (ldda < max(1,m))
*info = -5;
else if (lddb < max(1,m))
*info = -9;
else if (lwork < lwkopt && ! lquery)
*info = -11;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery)
return *info;
min_mn = min(m,n);
if (min_mn == 0) {
hwork[0] = c_one;
return *info;
}
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
/* B := Q^H * B */
magma_cunmqr_gpu( MagmaLeft, Magma_ConjTrans,
m, nrhs, n,
dA(0,0), ldda, tau,
dB, lddb, hwork, lwork, dT, nb, info );
if ( *info != 0 ) {
magma_queue_destroy( queue );
return *info;
}
/* Solve R*X = B(1:n,:) */
lddwork= min_mn;
if (nb < min_mn)
dwork = dT+2*lddwork*nb;
else
dwork = dT;
// To do: Why did we have this line originally; seems to be a bug (Stan)?
// dwork = dT;
i = (min_mn - 1)/nb * nb;
ib = n-i;
rows = m-i;
// TODO: this assumes that, on exit from magma_cunmqr_gpu, hwork contains
// the last block of A and B (i.e., C in cunmqr). This should be fixed.
// Seems this data should already be on the GPU, so could switch to
// magma_ctrsm and drop the csetmatrix.
if ( nrhs == 1 ) {
blasf77_ctrsv( MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&ib, hwork, &rows,
hwork+rows*ib, &ione);
} else {
blasf77_ctrsm( MagmaLeftStr, MagmaUpperStr, MagmaNoTransStr, MagmaNonUnitStr,
&ib, &nrhs,
&c_one, hwork, &rows,
hwork+rows*ib, &rows);
}
// update the solution vector
magma_csetmatrix( ib, nrhs,
hwork+rows*ib, rows,
dwork+i, lddwork, queue );
// update c
if (nrhs == 1) {
magma_cgemv( MagmaNoTrans, i, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, 1,
c_one, dB, 1, queue );
}
else {
magma_cgemm( MagmaNoTrans, MagmaNoTrans, i, nrhs, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, lddwork,
c_one, dB, lddb, queue );
}
magma_int_t start = i-nb;
if (nb < min_mn) {
for (i = start; i >= 0; i -= nb) {
ib = min(min_mn - i, nb);
rows = m - i;
if (i + ib < n) {
if (nrhs == 1) {
magma_cgemv( MagmaNoTrans, ib, ib,
c_one, dT(i), ib,
dB+i, 1,
c_zero, dwork+i, 1, queue );
magma_cgemv( MagmaNoTrans, i, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, 1,
c_one, dB, 1, queue );
}
else {
magma_cgemm( MagmaNoTrans, MagmaNoTrans, ib, nrhs, ib,
c_one, dT(i), ib,
dB+i, lddb,
c_zero, dwork+i, lddwork, queue );
magma_cgemm( MagmaNoTrans, MagmaNoTrans, i, nrhs, ib,
c_neg_one, dA(0, i), ldda,
dwork + i, lddwork,
c_one, dB, lddb, queue );
}
}
}
}
magma_ccopymatrix( n, nrhs,
dwork, lddwork,
dB, lddb, queue );
magma_queue_destroy( queue );
return *info;
}
extern "C" magma_int_t
magma_cpotrf3_mgpu(magma_int_t num_gpus, char uplo, magma_int_t m, magma_int_t n,
magma_int_t off_i, magma_int_t off_j, magma_int_t nb,
magmaFloatComplex *d_lA[], magma_int_t ldda,
magmaFloatComplex *d_lP[], magma_int_t lddp,
magmaFloatComplex *a, magma_int_t lda, magma_int_t h,
magma_queue_t stream[][3], magma_event_t event[][5],
magma_int_t *info )
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CPOTRF computes the Cholesky factorization of a complex Hermitian
positive definite matrix dA.
Auxiliary subroutine for cpotrf2_ooc. It is multiple gpu interface to compute
Cholesky of a "rectangular" matrix.
The factorization has the form
dA = U**H * U, if UPLO = 'U', or
dA = L * L**H, if UPLO = 'L',
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': Upper triangle of dA is stored;
= 'L': Lower triangle of dA is stored.
N (input) INTEGER
The order of the matrix dA. N >= 0.
dA (input/output) COMPLEX array on the GPU, dimension (LDDA,N)
On entry, the Hermitian matrix dA. If UPLO = 'U', the leading
N-by-N upper triangular part of dA contains the upper
triangular part of the matrix dA, and the strictly lower
triangular part of dA is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of dA contains the lower
triangular part of the matrix dA, and the strictly upper
triangular part of dA is not referenced.
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization dA = U**H * U or dA = L * L**H.
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
To benefit from coalescent memory accesses LDDA must be
dividable by 16.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
> 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
===================================================================== */
magma_int_t j, jb, nb0, nb2, d, dd, id, j_local, j_local2, buf;
char uplo_[2] = {uplo, 0};
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
float d_one = 1.0;
float d_neg_one = -1.0;
int upper = lapackf77_lsame(uplo_, "U");
magmaFloatComplex *dlpanel;
magma_int_t n_local[MagmaMaxGPUs], ldpanel;
const magma_int_t stream1 = 0, stream2 = 1, stream3 = 2;
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(CTRSM_WORK)
/* used by ctrsm_work */
int trsm_nb = 128;
int trsm_n = trsm_nb*((nb+trsm_nb-1)/trsm_nb);
magmaFloatComplex *d_dinvA[MagmaMaxGPUs];
magmaFloatComplex *d_x[MagmaMaxGPUs];
#define dinvA(d,j) &(d_dinvA[(d)][(j)*trsm_nb*trsm_n])
#define dx(d,j) &(d_x[(d)][(j)*nb*m])
/*
* Allocate device memory for the inversed diagonal blocks, size=N*BLOCK_SIZE
*/
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
if ( (MAGMA_SUCCESS != magma_cmalloc( &d_dinvA[d], 2*trsm_nb*trsm_n )) ||
(MAGMA_SUCCESS != magma_cmalloc( &d_x[d], 2*nb*(upper ? n : m) )) ) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
}
magma_setdevice(0);
#endif
*info = 0;
if ( (! upper) && (! lapackf77_lsame(uplo_, "L")) ) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (!upper && num_gpus*ldda < max(1,n)) {
*info = -4;
} else if (upper && ldda < max(1,m)) {
*info = -4;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* initialization */
for( d=0; d<num_gpus; d++ ) {
/* local-n and local-ld */
if (upper) {
n_local[d] = (n/(nb*num_gpus))*nb;
if (d < (n/nb)%num_gpus)
n_local[d] += nb;
else if (d == (n/nb)%num_gpus)
n_local[d] += n%nb;
} else {
n_local[d] = (m/(nb*num_gpus))*nb;
if (d < (m/nb)%num_gpus)
n_local[d] += nb;
else if (d == (m/nb)%num_gpus)
n_local[d] += m%nb;
}
}
/* == initialize the trace */
trace_init( 1, num_gpus, 3, (CUstream_st**)stream );
if (upper)
{
/* ---------------------------------------------- */
/* Upper-triangular case */
/* > Compute the Cholesky factorization A = U'*U. */
/* ---------------------------------------------- */
for (j=0; j<m; j+=nb) {
/* Set the GPU number that holds the current panel */
id = (j/nb)%num_gpus;
buf = (j/nb)%num_gpus; // right now, we have num_gpu buffers, so id and buf are the same..
/* Set the local index where the current panel is */
j_local = j/(nb*num_gpus);
jb = min(nb, (m-j));
/* Update the current diagonal block on stream1 */
magma_setdevice(id);
if( j > 0 ) {
magmablasSetKernelStream(stream[id][stream1]);
trace_gpu_start( id, stream1, "syrk", "syrk" );
magma_cherk(MagmaUpper, MagmaConjTrans, jb, j,
d_neg_one, dlA(id, 0, nb*j_local), ldda,
d_one, dlA(id, j, nb*j_local), ldda);
trace_gpu_end( id, stream1 );
}
/* send the diagonal to cpu on stream1 */
trace_gpu_start( id, stream1, "comm", "D to CPU" );
magma_cgetmatrix_async( jb, jb,
dlA(id, j, nb*j_local), ldda,
Aup(j,j), lda,
stream[id][stream1] );
trace_gpu_end( id, stream1 );
/* update off-diagonal blocks in the panel */
if( j > 0 ) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
j_local2 = j_local+1;
if( d > id ) j_local2 --;
nb0 = nb*j_local2; // number of local columns in the panel, while jb is panel-size (number of rows)
if( n_local[d] > nb0 ) {
magma_setdevice(d);
magmablasSetKernelStream(stream[d][stream2]);
if( d == id ) {
dlpanel = dlA(d,0,nb*j_local);
ldpanel = ldda;
// the GPU owns the row from start, and no need of synch.
//magma_queue_wait_event( stream[d][stream2], event[d][0] ); // rows arrived at gpu
} else {
dlpanel = dlP(d,nb,0,buf);
ldpanel = lddp;
magma_queue_wait_event( stream[d][stream2], event[d][0] ); // rows arrived at gpu
}
trace_gpu_start( d, stream2, "gemm", "gemm" );
magma_cgemm(MagmaConjTrans, MagmaNoTrans,
jb, n_local[d]-nb0, j,
c_neg_one, dlpanel, ldpanel,
dlA(d, 0, nb0), ldda,
c_one, dlA(d, j, nb0), ldda);
trace_gpu_end( d, stream2 );
magma_event_record( event[d][2], stream[d][stream2] );
}
d = (d+1)%num_gpus;
}
}
/* wait for panel and factorize it on cpu */
magma_setdevice(id);
magma_queue_sync( stream[id][stream1] );
trace_cpu_start( 0, "getrf", "getrf" );
lapackf77_cpotrf(MagmaUpperStr, &jb, Aup(j,j), &lda, info);
trace_cpu_end( 0 );
if (*info != 0) {
*info = *info + j;
break;
}
/* send the diagonal to gpus on stream1 */
if ( (j+jb) < n) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
if( d == id ) {
dlpanel = dlA(d, j, nb*j_local);
ldpanel = ldda;
} else {
dlpanel = dlP(d,0,0,buf);
ldpanel = lddp;
}
magma_setdevice(d);
trace_gpu_start( d, stream1, "comm", "comm" );
magma_csetmatrix_async( jb, jb,
Aup(j,j), lda,
dlpanel, ldpanel,
stream[d][stream1] );
trace_gpu_end( d, stream1 );
magma_event_record( event[d][1], stream[d][stream1] );
d = (d+1)%num_gpus;
}
} else {
magma_setdevice(id);
trace_gpu_start( id, stream1, "comm", "comm" );
magma_csetmatrix_async( jb, jb,
Aup(j,j), lda,
dlA(id, j, nb*j_local), ldda,
stream[id][stream1] );
trace_gpu_end( id, stream1 );
}
/* panel-factorize the off-diagonal */
if ( (j+jb) < n) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
/* next column */
j_local2 = j_local+1;
if( d > id ) j_local2--;
if( d == id ) {
dlpanel = dlA(d,j,nb*j_local);
ldpanel = ldda;
} else {
dlpanel = dlP(d,0,0,buf);
ldpanel = lddp;
}
nb2 = n_local[d] - j_local2*nb;
magma_setdevice(d);
if( j+jb < m && d == (j/nb+1)%num_gpus ) {
/* owns the next column, look-ahead next block on stream1 */
nb0 = min(nb, nb2);
magmablasSetKernelStream(stream[d][stream1]);
magma_queue_wait_event( stream[d][stream1], event[d][2] ); // wait for gemm update
trace_gpu_start( d, stream1, "trsm", "trsm" );
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(CTRSM_WORK)
magmablas_claset( MagmaUpperLower, trsm_nb, trsm_n, dinvA(d,0),trsm_nb );
magmablas_claset( MagmaUpperLower, nb0,jb, dx(d,0),nb0 );
magmablas_ctrsm_work( MagmaLeft, MagmaUpper,
MagmaConjTrans, MagmaNonUnit,
jb, nb0, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda,
1, dinvA(d,0), dx(d,0) );
#else
magma_ctrsm( MagmaLeft, MagmaUpper,
MagmaConjTrans, MagmaNonUnit,
jb, nb0, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda);
#endif
magma_event_record( event[d][4], stream[d][stream1] );
trace_gpu_end( d, stream1 );
} else if( nb2 > 0 ) {
/* update all the blocks on stream2 */
magma_queue_wait_event( stream[d][stream2], event[d][1] ); // wait for cholesky factor
trace_gpu_start( d, stream2, "trsm", "trsm" );
magmablasSetKernelStream(stream[d][stream2]);
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(CTRSM_WORK)
magmablas_claset( MagmaUpperLower, trsm_nb,trsm_n, dinvA(d,0),trsm_nb );
magmablas_claset( MagmaUpperLower, nb2,jb, dx(d,0),nb2 );
magmablas_ctrsm_work( MagmaLeft, MagmaUpper,
MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda,
1, dinvA(d,0), dx(d,0) );
#else
magma_ctrsm( MagmaLeft, MagmaUpper,
MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2), ldda);
#endif
trace_gpu_end( d, stream2 );
}
d = (d+1)%num_gpus;
} /* end of for */
/* ========================================================== */
if( j+jb < m ) {
d = (j/nb+1)%num_gpus;
/* next column */
j_local2 = j_local+1;
if( d > id ) j_local2--;
nb0 = min(nb, n_local[d]-nb*j_local2 );
/* even on 1 gpu, off-diagonals are copied to cpu (synchronize at the end). *
* so we have the Cholesky factor, but only diagonal submatrix of the big panel, *
* on cpu at the end. */
int d2, buf2;
magma_setdevice(d);
/* lookahead done */
magma_queue_wait_event( stream[d][stream3], event[d][4] );
trace_gpu_start( d, stream3, "comm", "row to CPU" );
magma_cgetmatrix_async( (j+jb), nb0,
dlA(d, 0, nb*j_local2), ldda,
Aup(0,j+jb), lda,
stream[d][stream3] );
trace_gpu_end( d, stream3 );
magma_event_record( event[d][3], stream[d][stream3] );
/* needed on pluto */
//magma_queue_sync( stream[d][stream3] );
/* broadcast rows to gpus on stream2 */
buf2 = ((j+jb)/nb)%num_gpus;
for( d2=0; d2<num_gpus; d2++ ) {
if( d2 != d )
{
magma_setdevice(d2);
trace_gpu_start( d2, stream3, "comm", "row to GPUs" );
magma_queue_wait_event( stream[d2][stream3], event[d][3] ); // rows arrived at cpu on stream3
magma_csetmatrix_async( j+jb, nb0,
Aup(0,j+jb), lda,
dlP(d2,nb,0,buf2), lddp,
stream[d2][stream3] );
trace_gpu_end( d2, stream3 );
magma_event_record( event[d2][0], stream[d2][stream3] );
}
}
/* =========================== */
/* update the remaining blocks */
nb2 = n_local[d]-(nb*j_local2 + nb0);
if( nb2 > 0 ) {
if( d == id ) {
dlpanel = dlA(d, j, nb*j_local);
ldpanel = ldda;
} else {
dlpanel = dlP(d,0,0,buf);
ldpanel = lddp;
}
magma_setdevice(d);
magmablasSetKernelStream(stream[d][stream2]);
trace_gpu_start( d, stream2, "trsm", "trsm" );
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(CTRSM_WORK)
int flag = 0;
if (flag == 0) {
magma_queue_wait_event( stream[d][stream2], event[d][4] ); // lookahead -> diagonal inversion
} else {
magmablas_claset( MagmaUpperLower, trsm_nb,trsm_n, dinvA(d,flag),trsm_nb );
magma_queue_wait_event( stream[d][stream2], event[d][1] ); // panel received
}
magmablas_claset( MagmaUpperLower, nb2,jb, dx(d,1),nb2 );
magmablas_ctrsm_work( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2+nb0), ldda,
flag, dinvA(d,flag), dx(d,1) );
#else
magma_queue_wait_event( stream[d][stream2], event[d][1] ); // wait for cholesky factor
magma_ctrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
jb, nb2, c_one,
dlpanel, ldpanel,
dlA(d, j, nb*j_local2+nb0), ldda);
#endif
trace_gpu_end( d, stream2 );
}
}
} /* end of ctrsm */
} /* end of for j=1, .., n */
} else {
/* ---------------------------------------------- */
/* Lower-triangular case */
/* > Compute the Cholesky factorization A = L*L'. */
/* ---------------------------------------------- */
for (j=0; j<n; j+=nb) {
/* Set the GPU number that holds the current panel */
id = (j/nb)%num_gpus;
buf = (j/nb)%num_gpus;
/* Set the local index where the current panel is */
j_local = j/(nb*num_gpus);
jb = min(nb, (n-j));
/* Update the current diagonal block on stream1 */
magma_setdevice(id);
if( j > 0 ) {
magmablasSetKernelStream(stream[id][stream1]);
magma_cherk(MagmaLower, MagmaNoTrans, jb, j,
d_neg_one, dlA(id, nb*j_local, 0), ldda,
d_one, dlA(id, nb*j_local, j), ldda);
}
/* send the diagonal to cpu on stream1 */
magma_cgetmatrix_async( jb, jb,
dlA(id, nb*j_local, j), ldda,
Alo(j,j), lda,
stream[id][stream1] );
/* update off-diagonal blocks of the panel */
if( j > 0 ) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
j_local2 = j_local+1;
if( d > id ) j_local2 --;
nb0 = nb*j_local2;
if( nb0 < n_local[d] ) {
magma_setdevice(d);
magmablasSetKernelStream(stream[d][stream2]);
if( d == id ) {
dlpanel = dlA(d, nb*j_local, 0);
ldpanel = ldda;
} else {
dlpanel = dlPT(d,0,nb,buf);
ldpanel = nb;
magma_queue_wait_event( stream[d][stream2], event[d][0] ); // rows arrived at gpu
}
magma_cgemm( MagmaNoTrans, MagmaConjTrans,
n_local[d]-nb0, jb, j,
c_neg_one, dlA(d, nb0, 0), ldda,
dlpanel, ldpanel,
c_one, dlA(d, nb0, j), ldda);
magma_event_record( event[d][2], stream[d][stream2] );
}
d = (d+1)%num_gpus;
}
}
/* wait for the panel and factorized it on cpu */
magma_setdevice(id);
magma_queue_sync( stream[id][stream1] );
lapackf77_cpotrf(MagmaLowerStr, &jb, Alo(j,j), &lda, info);
if (*info != 0) {
*info = *info + j;
break;
}
/* send the diagonal to gpus on stream1 */
if ( (j+jb) < m) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
if( d == id ) {
dlpanel = dlA(d, nb*j_local, j);
ldpanel = ldda;
} else {
dlpanel = dlPT(d, 0, 0, buf);
ldpanel = nb;
}
magma_setdevice(d);
magma_csetmatrix_async( jb, jb,
Alo(j,j), lda,
dlpanel, ldpanel,
stream[d][stream1] );
magma_event_record( event[d][1], stream[d][stream1] );
d = (d+1)%num_gpus;
}
} else {
magma_setdevice(id);
magma_csetmatrix_async( jb, jb,
Alo(j,j), lda,
dlA(id, nb*j_local, j), ldda,
stream[id][stream1] );
}
/* panel factorize the off-diagonal */
if ( (j+jb) < m) {
d = (j/nb+1)%num_gpus;
for( dd=0; dd<num_gpus; dd++ ) {
/* next column */
j_local2 = j_local+1;
if( d > id ) j_local2--;
if( d == id ) {
dlpanel = dlA(d, nb*j_local, j);
ldpanel = ldda;
} else {
dlpanel = dlPT(d, 0, 0, buf);
ldpanel = nb;
}
nb2 = n_local[d] - j_local2*nb;
nb0 = min(nb, nb2);
magma_setdevice(d);
if( j+nb < n && d == (j/nb+1)%num_gpus ) { /* owns next column, look-ahead next block on stream1 */
if ( j > 0 ) magma_queue_wait_event( stream[d][stream1], event[d][2] ); // wait for gemm update
magmablasSetKernelStream(stream[d][stream1]);
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(CTRSM_WORK)
magmablas_claset( MagmaUpperLower, trsm_nb, trsm_n, dinvA(d,0),trsm_nb );
magmablas_claset( MagmaUpperLower, nb0,jb, dx(d,0),nb0 );
magmablas_ctrsm_work( MagmaRight, MagmaLower,
MagmaConjTrans, MagmaNonUnit,
nb0, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda,
1, dinvA(d,0), dx(d,0) );
#else
magma_ctrsm( MagmaRight, MagmaLower,
MagmaConjTrans, MagmaNonUnit,
nb0, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda);
#endif
magma_event_record( event[d][4], stream[d][stream1] );
} else if( nb2 > 0 ) { /* other gpus updating all the blocks on stream2 */
/* update the entire column */
magma_queue_wait_event( stream[d][stream2], event[d][1] ); // wait for the cholesky factor
magmablasSetKernelStream(stream[d][stream2]);
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(CTRSM_WORK)
magmablas_claset( MagmaUpperLower, trsm_nb,trsm_n, dinvA(d,0),trsm_nb );
magmablas_claset( MagmaUpperLower, nb2,jb, dx(d,0),nb2 );
magmablas_ctrsm_work( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda,
1, dinvA(d,0), dx(d,0) );
#else
magma_ctrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2, j), ldda);
#endif
}
d = (d+1)%num_gpus;
} /* end for d */
/* ========================================================== */
if( j+jb < n ) {
d = (j/nb+1)%num_gpus;
/* next column */
j_local2 = j_local+1;
if( d > id ) j_local2--;
nb0 = min(nb, n_local[d]-nb*j_local2 );
/* even on 1 gpu, we copy off-diagonal to cpu (but don't synchronize). */
/* so we have the Cholesky factor on cpu at the end. */
int d2, buf2;
//#define CPOTRF_DEVICE_TO_DEVICE
#ifdef CPOTRF_DEVICE_TO_DEVICE
// lookahead done
/* broadcast the rows to gpus */
buf2 = ((j+jb)/nb)%num_gpus;
for( d2=0; d2<num_gpus; d2++ ) {
magma_setdevice(d2);
magma_queue_wait_event( stream[d2][stream3], event[d][4] );
if( d2 != d ) {
magma_ccopymatrix_async( nb0, j+jb,
dlPT(d2,0,nb,buf2), nb, // first nbxnb reserved for diagonal block
dlA(d, nb*j_local2, 0), ldda,
stream[d2][stream3] );
magma_event_record( event[d2][0], stream[d2][stream3] );
} else {
magma_cgetmatrix_async( nb0, j+jb,
dlA(d, nb*j_local2, 0), ldda,
Alo(j+jb,0), lda,
stream[d][stream3] );
}
}
#else
// lookahead done
magma_setdevice(d);
magma_queue_wait_event( stream[d][stream3], event[d][4] );
magma_cgetmatrix_async( nb0, j+jb,
dlA(d, nb*j_local2, 0), ldda,
Alo(j+jb,0), lda,
stream[d][stream3] );
magma_event_record( event[d][3], stream[d][stream3] );
/* syn on rows on CPU, seem to be needed on Pluto */
//magma_queue_sync( stream[d][stream3] );
/* broadcast the rows to gpus */
buf2 = ((j+jb)/nb)%num_gpus;
for( d2=0; d2<num_gpus; d2++ ) {
if( d2 != d )
{
magma_setdevice(d2);
magma_queue_wait_event( stream[d2][stream3], event[d][3] ); // getmatrix done
magma_csetmatrix_async( nb0, j+jb,
Alo(j+jb,0), lda,
dlPT(d2,0,nb,buf2), nb, // first nbxnb reserved for diagonal block
stream[d2][stream3] );
magma_event_record( event[d2][0], stream[d2][stream3] );
}
}
#endif
/* =================================== */
/* updates remaining blocks on stream2 */
nb2 = n_local[d] - (j_local2*nb + nb0);
if( nb2 > 0 ) {
if( d == id ) {
dlpanel = dlA(d, nb*j_local, j);
ldpanel = ldda;
} else {
dlpanel = dlPT(d,0,0,buf);
ldpanel = nb;
}
magma_setdevice(d);
magmablasSetKernelStream(stream[d][stream2]);
/* update the remaining blocks in the column */
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(CTRSM_WORK)
int flag = 0;
if (flag == 0) {
magma_queue_wait_event( stream[d][stream2], event[d][4] ); // lookahead -> diagonal inversion
} else {
magmablas_claset( MagmaUpperLower, trsm_nb,trsm_n, dinvA(d,flag),trsm_nb );
magma_queue_wait_event( stream[d][stream2], event[d][1] ); // panel received
}
magmablas_claset( MagmaUpperLower, nb2,jb, dx(d,1),nb2 );
magmablas_ctrsm_work( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2+nb0, j), ldda,
flag, dinvA(d,flag), dx(d,1) );
#else
magma_queue_wait_event( stream[d][stream2], event[d][1] ); // panel received
magma_ctrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
nb2, jb, c_one,
dlpanel, ldpanel,
dlA(d, nb*j_local2+nb0, j), ldda);
#endif
}
}
}
}
} /* end of else not upper */
/* == finalize the trace == */
trace_finalize( "cpotrf.svg","trace.css" );
for( d=0; d<num_gpus; d++ ) {
magma_setdevice(d);
for( j=0; j<3; j++ ) {
magma_queue_sync( stream[d][j] );
}
#if (defined(PRECISION_d) || defined(PRECISION_s)) && defined(CTRSM_WORK)
magma_free( d_dinvA[d] );
magma_free( d_x[d] );
#endif
magmablasSetKernelStream(NULL);
}
magma_setdevice(0);
return *info;
} /* magma_cpotrf_mgpu */
void magmablas_cher2k_mgpu2(
magma_uplo_t uplo, magma_trans_t trans, magma_int_t n, magma_int_t k,
magmaFloatComplex alpha,
magmaFloatComplex_ptr dA[], magma_int_t ldda, magma_int_t a_offset,
magmaFloatComplex_ptr dB[], magma_int_t lddb, magma_int_t b_offset,
float beta,
magmaFloatComplex_ptr dC[], magma_int_t lddc, magma_int_t c_offset,
magma_int_t ngpu, magma_int_t nb, magma_queue_t queues[][20], magma_int_t nqueue )
{
#define dA(dev, i, j) (dA[dev] + (i) + (j)*ldda + (a_offset) )
#define dB(dev, i, j) (dB[dev] + (i) + (j)*lddb + (b_offset) )
#define dC(dev, i, j) (dC[dev] + (i) + (j)*lddc)
/* Check arguments */
magma_int_t info = 0;
if ( uplo != MagmaLower ) {
info = -1; // upper not yet handled
} else if ( trans != MagmaNoTrans ) {
info = -2; // conj not yet handled
} else if ( n < 0 ) {
info = -3;
} else if ( k < 0 ) {
info = -4;
} else if ( ((trans == MagmaNoTrans) && ldda < max(1,n)) ||
((trans == Magma_ConjTrans) && ldda < max(1,k)) ) {
info = -7;
} else if ( a_offset < 0 || a_offset > ldda ) {
info = -8;
} else if ( ((trans == MagmaNoTrans) && lddb < max(1,n)) ||
((trans == Magma_ConjTrans) && lddb < max(1,k)) ) {
info = -10;
} else if ( b_offset < 0 || b_offset > lddb ) {
info = -11;
} else if ( lddc < max(1,n) ) {
info = -13;
} else if ( c_offset < 0 || c_offset > lddc ) {
info = -14;
} else if ( ngpu <= 0 ) {
info = -15;
} else if ( nb <= 0 ) {
info = -16;
} else if ( nqueue <= 0 ) {
info = -18;
}
if ( info != 0 ) {
magma_xerbla( __func__, -(info) );
return;
}
const magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex cbeta = MAGMA_C_MAKE( beta, 0. );
magma_int_t ib, ioff, iblock, idev, di, s;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
// loop over all blocks
// Faster to have two loops: first loop does C_hat = alpha*A*B**H + beta*C
// blockoffset is offset within first block; for subsequent blocks it is 0
magma_int_t blockoffset = c_offset % nb;
for( magma_int_t i = 0; i < n; i += ib ) {
ib = min( nb-blockoffset, n-i ); // block size
ioff = i + c_offset; // global index in parent matrix
iblock = (ioff / nb) / ngpu; // local block id
idev = (ioff / nb) % ngpu; // device with this block
di = iblock*nb + blockoffset; // local index in parent matrix
magma_setdevice( idev );
s = iblock % nqueue;
// C[i:n,i] = alpha * A[i:n,0] * B[i,0]' + beta*C[i:n,i]
//printf( "cgemm n=%4d, ib=%4d, k=%4d, i=%4d\n", n-i, ib, k, i );
magma_cgemm( MagmaNoTrans, Magma_ConjTrans, n-i, ib, k,
alpha, dA(idev,i,0), ldda,
dB(idev,i,0), lddb,
cbeta, dC(idev,ioff,di), lddc, queues[idev][s] );
blockoffset = 0;
}
// second loop does C = conj(alpha)*B*A**H + C_hat
alpha = MAGMA_C_CONJ( alpha );
blockoffset = c_offset % nb;
for( magma_int_t i = 0; i < n; i += ib ) {
ib = min( nb-blockoffset, n-i ); // block size
ioff = i + c_offset; // global index in parent matrix
iblock = (ioff / nb) / ngpu; // local block id
idev = (ioff / nb) % ngpu; // device with this block
di = iblock*nb + blockoffset; // local index in parent matrix
magma_setdevice( idev );
s = iblock % nqueue;
// C[i:n,i] += conj(alpha) * B[i:n,0] * A[i,0]'
//printf( "cgemm n=%4d, ib=%4d, k=%4d, i=%4d\n", n-i, ib, k, i );
magma_cgemm( MagmaNoTrans, Magma_ConjTrans, n-i, ib, k,
alpha, dB(idev,i,0), lddb,
dA(idev,i,0), ldda,
c_one, dC(idev,ioff,di), lddc, queues[idev][s] );
blockoffset = 0;
}
magma_setdevice( orig_dev );
}
/**
Purpose
-------
CLARFB applies a complex block reflector H or its transpose H^H to a
COMPLEX m by n matrix C, from the left.
__Note that this function assumes__ that the upper part of dV is 0
because it is referenced. Same for upper/lower part of dT.
Arguments
---------
@param[in]
side magma_side_t
- = MagmaLeft: apply H or H^H from the Left
- = MagmaRight: apply H or H^H from the Right
@param[in]
trans magma_trans_t
- = MagmaNoTrans: apply H (No transpose)
- = Magma_ConjTrans: apply H^H (Conjugate transpose)
@param[in]
direct magma_direct_t
Indicates how H is formed from a product of elementary
reflectors
- = MagmaForward: H = H(1) H(2) . . . H(k) (Forward)
- = MagmaBackward: H = H(k) . . . H(2) H(1) (Backward)
@param[in]
storev magma_storev_t
Indicates how the vectors which define the elementary
reflectors are stored:
- = MagmaColumnwise: Columnwise
- = MagmaRowwise: Rowwise
@param[in]
m INTEGER
The number of rows of the matrix C.
@param[in]
n INTEGER
The number of columns of the matrix C.
@param[in]
k INTEGER
The order of the matrix T (= the number of elementary
reflectors whose product defines the block reflector).
@param[in]
dV COMPLEX array on the GPU, dimension
(LDDV,K) if STOREV = MagmaColumnwise
(LDDV,M) if STOREV = MagmaRowwise and SIDE = MagmaLeft
(LDDV,N) if STOREV = MagmaRowwise and SIDE = MagmaRight
The matrix V. See further details.
@param[in]
lddv INTEGER
The leading dimension of the array V.
If STOREV = MagmaColumnwise and SIDE = MagmaLeft, LDDV >= max(1,M);
if STOREV = MagmaColumnwise and SIDE = MagmaRight, LDDV >= max(1,N);
if STOREV = MagmaRowwise, LDDV >= K.
@param[in]
dT COMPLEX array on the GPU, dimension (LDDT,K)
The triangular k by k matrix T in the representation of the
block reflector.
@param[in]
lddt INTEGER
The leading dimension of the array T. LDDT >= K.
@param[in,out]
dC COMPLEX array on the GPU, dimension (LDDC,N)
On entry, the m by n matrix C.
On exit, C is overwritten by H*C, or H^H*C, or C*H, or C*H^H.
@param[in]
lddc INTEGER
The leading dimension of the array C. LDA >= max(1,M).
@param
dwork (workspace) COMPLEX array, dimension (LDWORK,K)
@param[in]
ldwork INTEGER
The leading dimension of the array WORK.
If SIDE = MagmaLeft, LDWORK >= max(1,N);
if SIDE = MagmaRight, LDWORK >= max(1,M);
@param
dworkvt (workspace) COMPLEX array, dimension (LDWORKT,K)
@param[in]
ldworkvt INTEGER
The leading dimension of the array WORKVT.
LDWORKVT >= max(1,min(M,N));
Further Details
---------------
The shape of the matrix V and the storage of the vectors which define
the H(i) is best illustrated by the following example with n = 5 and
k = 3.
All elements including 0's and 1's are stored, unlike LAPACK.
DIRECT = MagmaForward and DIRECT = MagmaForward and
STOREV = MagmaColumnwise: STOREV = MagmaRowwise:
V = ( 1 0 0 ) V = ( 1 v1 v1 v1 v1 )
( v1 1 0 ) ( 0 1 v2 v2 v2 )
( v1 v2 1 ) ( 0 0 1 v3 v3 )
( v1 v2 v3 )
( v1 v2 v3 )
DIRECT = MagmaBackward and DIRECT = MagmaBackward and
STOREV = MagmaColumnwise: STOREV = MagmaRowwise:
V = ( v1 v2 v3 ) V = ( v1 v1 1 0 0 )
( v1 v2 v3 ) ( v2 v2 v2 1 0 )
( 1 v2 v3 ) ( v3 v3 v3 v3 1 )
( 0 1 v3 )
( 0 0 1 )
@ingroup magma_caux3
********************************************************************/
extern "C" magma_int_t
magma_clarfb_gpu_gemm(
magma_side_t side, magma_trans_t trans, magma_direct_t direct, magma_storev_t storev,
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloatComplex_const_ptr dV, magma_int_t lddv,
magmaFloatComplex_const_ptr dT, magma_int_t lddt,
magmaFloatComplex_ptr dC, magma_int_t lddc,
magmaFloatComplex_ptr dwork, magma_int_t ldwork,
magmaFloatComplex_ptr dworkvt, magma_int_t ldworkvt)
{
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t info = 0;
/* Function Body */
if (m <= 0 || n <= 0) {
return info;
}
// internal variable
magma_int_t ldwvt = (m > n ? k : m);
magma_int_t ldw;
if ( side == MagmaLeft ) {
ldw = k;
} else {
ldw = m;
}
// opposite of trans
magma_trans_t transt;
if (trans == MagmaNoTrans)
transt = Magma_ConjTrans;
else
transt = MagmaNoTrans;
MAGMA_UNUSED( transt ); // TODO: is this a bug that it isn't used?
// whether V is stored transposed or not
magma_trans_t notransV, transV;
if (storev == MagmaColumnwise) {
notransV = MagmaNoTrans;
transV = Magma_ConjTrans;
}
else {
notransV = Magma_ConjTrans;
transV = MagmaNoTrans;
}
if ( side == MagmaLeft ) {
// Form H C or H^H C
// Comments assume H C.
// When forming H^H C, T gets transposed via transt for m >= n or by trans for m < n.
// W = V^H C
magma_cgemm( Magma_ConjTrans, notransV,
k, n, m,
c_one, dV, lddv,
dC, lddc,
c_zero, dwork, ldw);
if (m <= n) {
// W2 = V T
magma_cgemm( notransV, trans,
m, k, k,
c_one, dV, lddv,
dT, lddt,
c_zero, dworkvt, ldwvt);
// C = C - W2 W = C - V T V^H C = (I - V T V^H) C = H C
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
m, n, k,
c_neg_one, dworkvt, ldwvt,
dwork, ldw,
c_one, dC, lddc);
} else {
// W2 = T W = T V^H C
magma_cgemm( trans, MagmaNoTrans,
k, n, k,
c_one, dT, lddt,
dwork, ldw,
c_zero, dworkvt, ldwvt);
// C = C - V W2 = C - V T V^H C = (I - V T V^H) C = H C
magma_cgemm( notransV, MagmaNoTrans,
m, n, k,
c_neg_one, dV, lddv,
dworkvt, ldwvt,
c_one, dC, lddc);
}
}
else {
// Form C H or C H^H
// Comments assume C H.
// When forming C H^H, T gets transposed via trans.
// W = C V
magma_cgemm( MagmaNoTrans, notransV,
m, k, n,
c_one, dC, lddc,
dV, lddv,
c_zero, dwork, ldw);
if (m <= n) {
// W2 = W T = C V T
magma_cgemm( MagmaNoTrans, trans,
m, k, k,
c_one, dwork, ldw,
dT, lddt,
c_zero, dworkvt, ldwvt);
// C = C - W2 V^H = C - C V T V^H = C (I - V T V^H) = C H
magma_cgemm( MagmaNoTrans, transV,
m, n, k,
c_neg_one, dworkvt, ldwvt,
dV, lddv,
c_one, dC, lddc);
} else {
// W2 = T V^H
magma_cgemm( trans, transV,
k, n, k,
c_one, dT, lddt,
dV, lddv,
c_zero, dworkvt, ldwvt);
// C = C - W W2 = C - C V T V^H = C (I - V T V^H) = C H
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
m, n, k,
c_neg_one, dwork, ldw,
dworkvt, ldwvt,
c_one, dC, lddc);
}
}
return info;
} /* magma_clarfb */
extern "C" magma_int_t
magma_cgetrf_gpu(magma_int_t m, magma_int_t n,
magmaFloatComplex *dA, magma_int_t ldda,
magma_int_t *ipiv, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
If the current stream is NULL, this version replaces it with user defined
stream to overlap computation with communication.
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.
A (input/output) COMPLEX array on the GPU, dimension (LDDA,N).
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
LDDA (input) INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
IPIV (output) INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
> 0: if INFO = i, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
===================================================================== */
#define dAT(i,j) (dAT + (i)*nb*lddat + (j)*nb)
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t iinfo, nb;
magma_int_t maxm, maxn, mindim;
magma_int_t i, rows, cols, s, lddat, lddwork;
magmaFloatComplex *dAT, *dAP, *work;
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
mindim = min(m, n);
nb = magma_get_cgetrf_nb(m);
s = mindim / nb;
if (nb <= 1 || nb >= min(m,n)) {
/* Use CPU code. */
magma_cmalloc_cpu( &work, m * n );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
magma_cgetmatrix( m, n, dA, ldda, work, m );
lapackf77_cgetrf(&m, &n, work, &m, ipiv, info);
magma_csetmatrix( m, n, work, m, dA, ldda );
magma_free_cpu(work);
}
else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
maxn = ((n + 31)/32)*32;
lddat = maxn;
lddwork = maxm;
dAT = dA;
if (MAGMA_SUCCESS != magma_cmalloc( &dAP, nb*maxm )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
if ( m == n ) {
lddat = ldda;
magmablas_ctranspose_inplace( m, dAT, ldda );
}
else {
if (MAGMA_SUCCESS != magma_cmalloc( &dAT, maxm*maxn )) {
magma_free( dAP );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magmablas_ctranspose2( dAT, lddat, dA, ldda, m, n );
}
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &work, maxm*nb )) {
magma_free( dAP );
if ( ! (m == n))
magma_free( dAT );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
/* Define user stream if current stream is NULL */
cudaStream_t stream[2], current_stream;
magmablasGetKernelStream(¤t_stream);
magma_queue_create( &stream[0] );
if (current_stream == NULL) {
magma_queue_create( &stream[1] );
magmablasSetKernelStream(stream[1]);
}
else
stream[1] = current_stream;
for( i=0; i<s; i++ )
{
// download i-th panel
cols = maxm - i*nb;
//magmablas_ctranspose( dAP, cols, dAT(i,i), lddat, nb, cols );
magmablas_ctranspose2( dAP, cols, dAT(i,i), lddat, nb, m-i*nb );
// make sure that that the transpose has completed
magma_queue_sync( stream[1] );
magma_cgetmatrix_async( m-i*nb, nb, dAP, cols, work, lddwork,
stream[0]);
if ( i>0 ){
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (i+1)*nb, nb,
c_one, dAT(i-1,i-1), lddat,
dAT(i-1,i+1), lddat );
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(i+1)*nb, m-i*nb, nb,
c_neg_one, dAT(i-1,i+1), lddat,
dAT(i, i-1), lddat,
c_one, dAT(i, i+1), lddat );
}
// do the cpu part
rows = m - i*nb;
magma_queue_sync( stream[0] );
lapackf77_cgetrf( &rows, &nb, work, &lddwork, ipiv+i*nb, &iinfo);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + i*nb;
// upload i-th panel
magma_csetmatrix_async( m-i*nb, nb, work, lddwork, dAP, maxm,
stream[0]);
magmablas_cpermute_long2( n, dAT, lddat, ipiv, nb, i*nb );
magma_queue_sync( stream[0] );
//magmablas_ctranspose(dAT(i,i), lddat, dAP, maxm, cols, nb);
magmablas_ctranspose2(dAT(i,i), lddat, dAP, maxm, m-i*nb, nb);
// do the small non-parallel computations (next panel update)
if ( s > (i+1) ) {
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dAT(i, i ), lddat,
dAT(i, i+1), lddat);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
nb, m-(i+1)*nb, nb,
c_neg_one, dAT(i, i+1), lddat,
dAT(i+1, i ), lddat,
c_one, dAT(i+1, i+1), lddat );
}
else {
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb, nb,
c_one, dAT(i, i ), lddat,
dAT(i, i+1), lddat);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(i+1)*nb, m-(i+1)*nb, nb,
c_neg_one, dAT(i, i+1), lddat,
dAT(i+1, i ), lddat,
c_one, dAT(i+1, i+1), lddat );
}
}
magma_int_t nb0 = min(m - s*nb, n - s*nb);
rows = m - s*nb;
cols = maxm - s*nb;
magmablas_ctranspose2( dAP, maxm, dAT(s,s), lddat, nb0, rows);
magma_cgetmatrix( rows, nb0, dAP, maxm, work, lddwork );
// do the cpu part
lapackf77_cgetrf( &rows, &nb0, work, &lddwork, ipiv+s*nb, &iinfo);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + s*nb;
magmablas_cpermute_long2( n, dAT, lddat, ipiv, nb0, s*nb );
// upload i-th panel
magma_csetmatrix( rows, nb0, work, lddwork, dAP, maxm );
magmablas_ctranspose2( dAT(s,s), lddat, dAP, maxm, rows, nb0);
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb-nb0, nb0,
c_one, dAT(s,s), lddat,
dAT(s,s)+nb0, lddat);
if ( m == n ) {
magmablas_ctranspose_inplace( m, dAT, lddat );
}
else {
magmablas_ctranspose2( dA, ldda, dAT, lddat, n, m );
magma_free( dAT );
}
magma_free( dAP );
magma_free_pinned( work );
magma_queue_destroy( stream[0] );
if (current_stream == NULL) {
magma_queue_destroy( stream[1] );
magmablasSetKernelStream(NULL);
}
}
return *info;
} /* End of MAGMA_CGETRF_GPU */
/**
@deprecated
Purpose
-------
CLAQPS 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]
A COMPLEX 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.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= 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 array, dimension (KB)
The scalar factors of the elementary reflectors.
@param[in,out]
vn1 REAL array, dimension (N)
The vector with the partial column norms.
@param[in,out]
vn2 REAL array, dimension (N)
The vector with the exact column norms.
@param[in,out]
auxv COMPLEX array, dimension (NB)
Auxiliar vector.
@param[in,out]
F COMPLEX array, dimension (LDF,NB)
Matrix F' = L*Y'*A.
@param[in]
ldf INTEGER
The leading dimension of the array F. LDF >= max(1,N).
@ingroup magma_cgeqp3_aux
********************************************************************/
extern "C" magma_int_t
magma_claqps_gpu(magma_int_t m, magma_int_t n, magma_int_t offset,
magma_int_t nb, magma_int_t *kb,
magmaFloatComplex *A, magma_int_t lda,
magma_int_t *jpvt, magmaFloatComplex *tau,
float *vn1, float *vn2,
magmaFloatComplex *auxv,
magmaFloatComplex *F, magma_int_t ldf)
{
#define A(i, j) (A + (i) + (j)*(lda ))
#define F(i, j) (F + (i) + (j)*(ldf ))
magmaFloatComplex c_zero = MAGMA_C_MAKE( 0.,0.);
magmaFloatComplex c_one = MAGMA_C_MAKE( 1.,0.);
magmaFloatComplex c_neg_one = MAGMA_C_MAKE(-1.,0.);
magma_int_t ione = 1;
magma_int_t i__1, i__2;
//float d__1;
magmaFloatComplex z__1;
//magma_int_t j;
magma_int_t k, rk;
//magmaFloatComplex Akk;
magmaFloatComplex *Aks;
magmaFloatComplex tauk = MAGMA_C_ZERO;
magma_int_t pvt;
//float temp, temp2;
float tol3z;
magma_int_t itemp;
float lsticc, *lsticcs;
magma_int_t lastrk;
magma_smalloc( &lsticcs, 1+256*(n+255)/256 );
lastrk = min( m, n + offset );
tol3z = magma_ssqrt( lapackf77_slamch("Epsilon"));
lsticc = 0;
k = 0;
magma_cmalloc( &Aks, nb );
while( k < nb && lsticc == 0 ) {
rk = offset + k;
/* Determine ith pivot column and swap if necessary */
// subtract 1 from Fortran/CUBLAS isamax; pvt, k are 0-based.
pvt = k + magma_isamax( n-k, &vn1[k], ione ) - 1;
if (pvt != k) {
/*if (pvt >= nb) {
// 1. Start copy from GPU
magma_cgetmatrix_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_cswap( &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_cswap(&m, A(0, pvt), &ione, A(0, k), &ione);
// 3. Restore the GPU
magma_csetmatrix_async( m - offset - nb, 1,
A (offset + nb, pvt), lda,
dA(offset + nb, pvt), ldda, stream);
}*/
magmablas_cswap( m, A(0, pvt), ione, A(0, k), ione );
//blasf77_cswap( &i__1, F(pvt,0), &ldf, F(k,0), &ldf );
magmablas_cswap( 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_C_CNJG( *F(k,j) );
}
#endif*/
//#define RIGHT_UPDATE
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_cgemv( 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_cgemv( MagmaNoTransStr, &i__1, &i__2,
&c_neg_one, A(rk, 0), &lda,
F(k, 0), &ldf,
&c_one, A(rk, k), &ione );*/
magma_cgemv( 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_C_CNJG( *F(k,j) );
}
#endif*/
}
/* Generate elementary reflector H(k). */
magma_clarfg_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_cgetvector( 1, &Aks[k], 1, &Akk, 1 );
/* needed to avoid the race condition */
if (k == 0) magma_csetvector( 1, &c_one, 1, A(rk, k), 1 );
else magma_ccopymatrix( 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_cgetvector( 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_csetmatrix( i__1, 1, A(rk, k), lda, dA(rk,k), ldda );
/* Multiply on GPU */
// was CALL CGEMV( '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_cgetvector( 1, &tau[k], 1, &tauk, 1 );
magma_cgemv( MagmaConjTrans, m-rk, n-k-1,
tauk, A( rk, k+1 ), lda,
A( rk, k ), 1,
c_zero, F( k+1, k ), 1 );
//magma_cscal( 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_cgemv( 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_cgetmatrix_async( i__2-i__3, 1,
// dF(k + 1 +i__3, k), i__2,
// F (k + 1 +i__3, k), i__2, stream );
//blasf77_cgemv( 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_cgemv( 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_csetvector( 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_cgetvector( 1, &tau[k], 1, &tauk, 1 );
z__1 = MAGMA_C_NEGATE( tauk );
#ifdef RIGHT_UPDATE
i__1 = m - offset - nb;
i__2 = k;
magma_cgemv( 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_cgemv( 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_cgemv( MagmaConjTransStr, &i__1, &i__2,
// &z__1, A(rk, 0), &lda,
// A(rk, k), &ione,
// &c_zero, auxv, &ione );
magma_cgemv( MagmaConjTrans, i__1, i__2,
z__1, A(rk, 0), lda,
A(rk, k), ione,
c_zero, auxv, ione );
//i__1 = k;
//blasf77_cgemv( MagmaNoTransStr, &n, &i__1,
// &c_one, F(0,0), &ldf,
// auxv, &ione,
// &c_one, F(0,k), &ione );
/*magma_cgemv( 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_cgemv( 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_cgemm( 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_cgemm( 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_cgemm( 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_scnrm2_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_sgetvector( 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_C_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] = (float) lsticc;
lsticc = j;
} else {
vn1[j] *= magma_ssqrt(temp);
}
}
}
}*/
//*A(rk, k) = Akk;
//magma_csetvector( 1, &Akk, 1, A(rk, k), 1 );
//magma_cswap( 1, &Aks[k], 1, A(rk, k), 1 );
++k;
}
magma_ccopymatrix( 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_csetmatrix( i__2, *kb,
F (*kb, 0), ldf,
dF(*kb, 0), i__2 );*/
magma_cgemm( 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_scnrm2_check(m-rk-1, n-*kb, A(rk+1,*kb), lda,
&vn1[*kb], lsticcs);
magma_scopymatrix( 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_scnrm2( i__1, A(rk+1,lsticc), ione );
else {
// Where is the data, CPU or GPU ?
float r1, r2;
r1 = magma_cblas_scnrm2( nb-k, A(rk+1,lsticc), ione );
r2 = magma_scnrm2(m-offset-nb, dA(offset + nb + 1, lsticc), ione);
vn1[lsticc] = magma_ssqrt(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(SLAMCH('S'))
vn2[lsticc] = vn1[lsticc];
lsticc = itemp;*/
}
magma_free(Aks);
magma_free(lsticcs);
return MAGMA_SUCCESS;
} /* magma_claqps */
/**
Purpose
-------
CGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges. This version does not
require work space on the GPU passed as input. GPU memory is allocated
in the routine.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
If the current stream is NULL, this version replaces it with user defined
stream to overlap computation with communication.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
A COMPLEX array, dimension (LDA,N)
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
\n
Higher performance is achieved if A is in pinned memory, e.g.
allocated using magma_malloc_pinned.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@param[out]
ipiv INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
- > 0: if INFO = i, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
@ingroup magma_cgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_cgetrf(magma_int_t m, magma_int_t n, magmaFloatComplex *A, magma_int_t lda,
magma_int_t *ipiv, magma_int_t *info)
{
#define dAT(i,j) (dAT + (i)*nb*ldda + (j)*nb)
magmaFloatComplex *dAT, *dA, *da, *work;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t iinfo, nb;
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (lda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
nb = magma_get_cgetrf_nb(m);
if ( (nb <= 1) || (nb >= min(m,n)) ) {
/* Use CPU code. */
lapackf77_cgetrf(&m, &n, A, &lda, ipiv, info);
} else {
/* Use hybrid blocked code. */
magma_int_t maxm, maxn, ldda, maxdim;
magma_int_t i, rows, cols, s = min(m, n)/nb;
maxm = ((m + 31)/32)*32;
maxn = ((n + 31)/32)*32;
maxdim = max(maxm, maxn);
/* set number of GPUs */
magma_int_t num_gpus = magma_num_gpus();
if ( num_gpus > 1 ) {
/* call multi-GPU non-GPU-resident interface */
magma_cgetrf_m(num_gpus, m, n, A, lda, ipiv, info);
return *info;
}
/* explicitly checking the memory requirement */
size_t freeMem, totalMem;
cudaMemGetInfo( &freeMem, &totalMem );
freeMem /= sizeof(magmaFloatComplex);
int h = 1+(2+num_gpus), num_gpus2 = num_gpus;
int NB = (magma_int_t)(0.8*freeMem/maxm-h*nb);
const char* ngr_nb_char = getenv("MAGMA_NGR_NB");
if ( ngr_nb_char != NULL )
NB = max( nb, min( NB, atoi(ngr_nb_char) ) );
if ( num_gpus > ceil((float)NB/nb) ) {
num_gpus2 = (int)ceil((float)NB/nb);
h = 1+(2+num_gpus2);
NB = (magma_int_t)(0.8*freeMem/maxm-h*nb);
}
if ( num_gpus2*NB < n ) {
/* require too much memory, so call non-GPU-resident version */
magma_cgetrf_m(num_gpus, m, n, A, lda, ipiv, info);
return *info;
}
ldda = maxn;
work = A;
if (maxdim*maxdim < 2*maxm*maxn) {
// if close to square, allocate square matrix and transpose in-place
if (MAGMA_SUCCESS != magma_cmalloc( &dA, nb*maxm + maxdim*maxdim )) {
/* alloc failed so call non-GPU-resident version */
magma_cgetrf_m(num_gpus, m, n, A, lda, ipiv, info);
return *info;
}
da = dA + nb*maxm;
ldda = maxdim;
magma_csetmatrix( m, n, A, lda, da, ldda );
dAT = da;
magmablas_ctranspose_inplace( ldda, dAT, ldda );
}
else {
// if very rectangular, allocate dA and dAT and transpose out-of-place
if (MAGMA_SUCCESS != magma_cmalloc( &dA, (nb + maxn)*maxm )) {
/* alloc failed so call non-GPU-resident version */
magma_cgetrf_m(num_gpus, m, n, A, lda, ipiv, info);
return *info;
}
da = dA + nb*maxm;
magma_csetmatrix( m, n, A, lda, da, maxm );
if (MAGMA_SUCCESS != magma_cmalloc( &dAT, maxm*maxn )) {
/* alloc failed so call non-GPU-resident version */
magma_free( dA );
magma_cgetrf_m(num_gpus, m, n, A, lda, ipiv, info);
return *info;
}
magmablas_ctranspose( m, n, da, maxm, dAT, ldda );
}
lapackf77_cgetrf( &m, &nb, work, &lda, ipiv, &iinfo);
/* Define user stream if current stream is NULL */
cudaStream_t stream[2], current_stream;
magmablasGetKernelStream(¤t_stream);
magma_queue_create( &stream[0] );
if (current_stream == NULL) {
magma_queue_create( &stream[1] );
magmablasSetKernelStream(stream[1]);
}
else
stream[1] = current_stream;
for( i = 0; i < s; i++ ) {
// download i-th panel
cols = maxm - i*nb;
if (i > 0) {
// download i-th panel
magmablas_ctranspose( nb, cols, dAT(i,i), ldda, dA, cols );
// make sure that gpu queue is empty
magma_device_sync();
magma_cgetmatrix_async( m-i*nb, nb, dA, cols, work, lda,
stream[0]);
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (i+1)*nb, nb,
c_one, dAT(i-1,i-1), ldda,
dAT(i-1,i+1), ldda );
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(i+1)*nb, m-i*nb, nb,
c_neg_one, dAT(i-1,i+1), ldda,
dAT(i, i-1), ldda,
c_one, dAT(i, i+1), ldda );
// do the cpu part
rows = m - i*nb;
magma_queue_sync( stream[0] );
lapackf77_cgetrf( &rows, &nb, work, &lda, ipiv+i*nb, &iinfo);
}
if (*info == 0 && iinfo > 0)
*info = iinfo + i*nb;
// upload i-th panel
magma_csetmatrix_async( m-i*nb, nb, work, lda, dA, cols,
stream[0]);
magmablas_cpermute_long2( ldda, dAT, ldda, ipiv, nb, i*nb );
magma_queue_sync( stream[0] );
magmablas_ctranspose( cols, nb, dA, cols, dAT(i,i), ldda );
// do the small non-parallel computations
if (s > (i+1)) {
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
nb, nb,
c_one, dAT(i, i ), ldda,
dAT(i, i+1), ldda);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
nb, m-(i+1)*nb, nb,
c_neg_one, dAT(i, i+1), ldda,
dAT(i+1, i ), ldda,
c_one, dAT(i+1, i+1), ldda );
}
else {
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb, nb,
c_one, dAT(i, i ), ldda,
dAT(i, i+1), ldda);
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(i+1)*nb, m-(i+1)*nb, nb,
c_neg_one, dAT(i, i+1), ldda,
dAT(i+1, i ), ldda,
c_one, dAT(i+1, i+1), ldda );
}
}
magma_int_t nb0 = min(m - s*nb, n - s*nb);
if ( nb0 > 0 ) {
rows = m - s*nb;
cols = maxm - s*nb;
magmablas_ctranspose( nb0, rows, dAT(s,s), ldda, dA, cols );
magma_cgetmatrix( rows, nb0, dA, cols, work, lda );
// make sure that gpu queue is empty
magma_device_sync();
// do the cpu part
lapackf77_cgetrf( &rows, &nb0, work, &lda, ipiv+s*nb, &iinfo);
if (*info == 0 && iinfo > 0)
*info = iinfo + s*nb;
magmablas_cpermute_long2( ldda, dAT, ldda, ipiv, nb0, s*nb );
magma_csetmatrix( rows, nb0, work, lda, dA, cols );
magmablas_ctranspose( rows, nb0, dA, cols, dAT(s,s), ldda );
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-s*nb-nb0, nb0,
c_one, dAT(s, s), ldda,
dAT(s, s)+nb0, ldda);
}
if (maxdim*maxdim < 2*maxm*maxn) {
magmablas_ctranspose_inplace( ldda, dAT, ldda );
magma_cgetmatrix( m, n, da, ldda, A, lda );
} else {
magmablas_ctranspose( n, m, dAT, ldda, da, maxm );
magma_cgetmatrix( m, n, da, maxm, A, lda );
magma_free( dAT );
}
magma_free( dA );
magma_queue_destroy( stream[0] );
if (current_stream == NULL) {
magma_queue_destroy( stream[1] );
magmablasSetKernelStream(NULL);
}
}
return *info;
} /* magma_cgetrf */
/***************************************************************************//**
Purpose
-------
CLAHRU is an auxiliary MAGMA routine that is used in CGEHRD to update
the trailing sub-matrices after the reductions of the corresponding
panels.
See further details below.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in]
ihi INTEGER
Last row to update. Same as IHI in cgehrd.
@param[in]
k INTEGER
Number of rows of the matrix Am (see details below)
@param[in]
nb INTEGER
Block size
@param[out]
A COMPLEX array, dimension (LDA,N-K)
On entry, the N-by-(N-K) general matrix to be updated. The
computation is done on the GPU. After Am is updated on the GPU
only Am(1:NB) is transferred to the CPU - to update the
corresponding Am matrix. See Further Details below.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[in,out]
dA COMPLEX array on the GPU, dimension (LDDA,N-K).
On entry, the N-by-(N-K) general matrix to be updated.
On exit, the 1st K rows (matrix Am) of A are updated by
applying an orthogonal transformation from the right
Am = Am (I-V T V'), and sub-matrix Ag is updated by
Ag = (I - V T V') Ag (I - V T V(NB+1:)' )
where Q = I - V T V' represent the orthogonal matrix
(as a product of elementary reflectors V) used to reduce
the current panel of A to upper Hessenberg form. After Am
is updated Am(:,1:NB) is sent to the CPU.
See Further Details below.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
@param[in,out]
dY (workspace) COMPLEX array on the GPU, dimension (LDDY, NB).
On entry the (N-K)-by-NB Y = A V. It is used internally
as workspace, so its value is changed on exit.
@param[in]
lddy INTEGER
The leading dimension of the array dY. LDDY >= max(1,N).
@param[in,out]
dV (workspace) COMPLEX array on the GPU, dimension (LDDV, NB).
On entry the (N-K)-by-NB matrix V of elementary reflectors
used to reduce the current panel of A to upper Hessenberg form.
The rest K-by-NB part is used as workspace. V is unchanged on
exit.
@param[in]
lddv INTEGER
The leading dimension of the array dV. LDDV >= max(1,N).
@param[in]
dT COMPLEX array on the GPU, dimension (NB, NB).
On entry the NB-by-NB upper trinagular matrix defining the
orthogonal Hessenberg reduction transformation matrix for
the current panel. The lower triangular part are 0s.
@param
dwork (workspace) COMPLEX array on the GPU, dimension N*NB.
@param[in]
queue magma_queue_t
Queue to execute in.
Further Details
---------------
This implementation follows the algorithm and notations described in:
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
The difference is that here Am is computed on the GPU.
M is renamed Am, G is renamed Ag.
@ingroup magma_lahru
*******************************************************************************/
extern "C" magma_int_t
magma_clahru(
magma_int_t n, magma_int_t ihi, magma_int_t k, magma_int_t nb,
magmaFloatComplex *A, magma_int_t lda,
magmaFloatComplex_ptr dA, magma_int_t ldda,
magmaFloatComplex_ptr dY, magma_int_t lddy,
magmaFloatComplex_ptr dV, magma_int_t lddv,
magmaFloatComplex_ptr dT,
magmaFloatComplex_ptr dwork,
magma_queue_t queue )
{
#define dA(i_,j_) (dA + (i_) + (j_)*ldda)
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *dYm = dV + ihi - k;
magma_int_t info = 0;
if (n < 0) {
info = -1;
} else if (ihi < 0 || ihi > n) {
info = -2;
} else if (k < 0 || k > n) {
info = -3;
} else if (nb < 1 || nb > n) {
info = -4;
} else if (lda < max(1,n)) {
info = -6;
} else if (ldda < max(1,n)) {
info = -8;
} else if (lddy < max(1,n)) {
info = -10;
} else if (lddv < max(1,n)) {
info = -12;
}
if (info != 0) {
magma_xerbla( __func__, -(info) );
return info;
}
// top part of Y, above panel, hasn't been computed yet, so do that now
// Ym = Am V = A(0:k-1, 0:ihi-k-1) * V(0:ihi-k-1, 0:nb-1)
magma_cgemm( MagmaNoTrans, MagmaNoTrans, k, nb, ihi-k,
c_one, dA, ldda,
dV, lddv,
c_zero, dYm, ldda, queue );
// -----
// on right, A := A Q = A - A V T V'
// Update Am = Am - Am V T V' = Am - Ym W', with W = V T'
// W = V T' = V(0:ihi-k-1, 0:nb-1) * T(0:nb-1, 0:nb-1)'
magma_cgemm( MagmaNoTrans, MagmaConjTrans, ihi-k, nb, nb,
c_one, dV, lddv,
dT, nb,
c_zero, dwork, ldda, queue );
// Am = Am - Ym W' = A(0:k-1, 0:ihi-k-1) - Ym(0:k-1, 0:nb-1) * W(0:ihi-k-1, 0:nb-1)'
magma_cgemm( MagmaNoTrans, MagmaConjTrans, k, ihi-k, nb,
c_neg_one, dYm, ldda,
dwork, ldda,
c_one, dA, ldda, queue );
// copy first nb columns of Am, A(0:k-1, 0:nb-1), to host
magma_cgetmatrix( k, nb, dA, ldda, A, lda, queue );
// -----
// on right, A := A Q = A - A V T V'
// Update Ag = Ag - Ag V T V' = Ag - Y W'
// Ag = Ag - Y W' = A(k:ihi-1, nb:ihi-k-1) - Y(0:ihi-k-1, 0:nb-1) * W(nb:ihi-k-1, 0:nb-1)'
magma_cgemm( MagmaNoTrans, MagmaConjTrans, ihi-k, ihi-k-nb, nb,
c_neg_one, dY, ldda,
dwork + nb, ldda,
c_one, dA(k,nb), ldda, queue );
// -----
// on left, A := Q' A = A - V T' V' A
// Ag2 = Ag2 - V T' V' Ag2 = W Yg, with W = V T' and Yg = V' Ag2
// Note that Ag is A(k:ihi, nb+1:ihi-k)
// while Ag2 is A(k:ihi, nb+1: n -k)
// Z = V(0:ihi-k-1, 0:nb-1)' * A(k:ihi-1, nb:n-k-1); Z is stored over Y
magma_cgemm( MagmaConjTrans, MagmaNoTrans, nb, n-k-nb, ihi-k,
c_one, dV, lddv,
dA(k,nb), ldda,
c_zero, dY, nb, queue );
// Ag2 = Ag2 - W Z = A(k:ihi-1, nb:n-k-1) - W(nb:n-k-1, 0:nb-1) * Z(0:nb-1, nb:n-k-1)
magma_cgemm( MagmaNoTrans, MagmaNoTrans, ihi-k, n-k-nb, nb,
c_neg_one, dwork, ldda,
dY, nb,
c_one, dA(k,nb), ldda, queue );
return info;
}
/**
Purpose
-------
CLARFB applies a complex block reflector H or its transpose H^H to a
COMPLEX m by n matrix C, from the left.
Arguments
---------
@param[in]
side magma_side_t
- = MagmaLeft: apply H or H^H from the Left
- = MagmaRight: apply H or H^H from the Right
@param[in]
trans magma_trans_t
- = MagmaNoTrans: apply H (No transpose)
- = Magma_ConjTrans: apply H^H (Conjugate transpose)
@param[in]
direct magma_direct_t
Indicates how H is formed from a product of elementary
reflectors
- = MagmaForward: H = H(1) H(2) . . . H(k) (Forward)
- = MagmaBackward: H = H(k) . . . H(2) H(1) (Backward)
@param[in]
storev magma_storev_t
Indicates how the vectors which define the elementary
reflectors are stored:
- = MagmaColumnwise: Columnwise
- = MagmaRowwise: Rowwise
@param[in]
m INTEGER
The number of rows of the matrix C.
@param[in]
n INTEGER
The number of columns of the matrix C.
@param[in]
k INTEGER
The order of the matrix T (= the number of elementary
reflectors whose product defines the block reflector).
@param[in]
dV COMPLEX array on the GPU, dimension
(LDV,K) if STOREV = MagmaColumnwise
(LDV,M) if STOREV = MagmaRowwise and SIDE = MagmaLeft
(LDV,N) if STOREV = MagmaRowwise and SIDE = MagmaRight
The matrix V. See further details.
@param[in]
ldv INTEGER
The leading dimension of the array V.
If STOREV = MagmaColumnwise and SIDE = MagmaLeft, LDV >= max(1,M);
if STOREV = MagmaColumnwise and SIDE = MagmaRight, LDV >= max(1,N);
if STOREV = MagmaRowwise, LDV >= K.
@param[in]
dT COMPLEX array on the GPU, dimension (LDT,K)
The triangular k by k matrix T in the representation of the
block reflector.
@param[in]
ldt INTEGER
The leading dimension of the array T. LDT >= K.
@param[in,out]
dC COMPLEX array on the GPU, dimension (LDC,N)
On entry, the m by n matrix C.
On exit, C is overwritten by H*C, or H^H*C, or C*H, or C*H^H.
@param[in]
ldc INTEGER
The leading dimension of the array C. LDA >= max(1,M).
@param
dwork (workspace) COMPLEX array, dimension (LDWORK,K)
@param[in]
ldwork INTEGER
The leading dimension of the array WORK.
If SIDE = MagmaLeft, LDWORK >= max(1,N);
if SIDE = MagmaRight, LDWORK >= max(1,M);
Further Details
---------------
The shape of the matrix V and the storage of the vectors which define
the H(i) is best illustrated by the following example with n = 5 and
k = 3.
All elements including 0's and 1's are stored, unlike LAPACK.
DIRECT = MagmaForward and DIRECT = MagmaForward and
STOREV = MagmaColumnwise: STOREV = MagmaRowwise:
V = ( 1 0 0 ) V = ( 1 v1 v1 v1 v1 )
( v1 1 0 ) ( 0 1 v2 v2 v2 )
( v1 v2 1 ) ( 0 0 1 v3 v3 )
( v1 v2 v3 )
( v1 v2 v3 )
DIRECT = MagmaBackward and DIRECT = MagmaBackward and
STOREV = MagmaColumnwise: STOREV = MagmaRowwise:
V = ( v1 v2 v3 ) V = ( v1 v1 1 0 0 )
( v1 v2 v3 ) ( v2 v2 v2 1 0 )
( 1 v2 v3 ) ( v3 v3 v3 v3 1 )
( 0 1 v3 )
( 0 0 1 )
@ingroup magma_caux3
********************************************************************/
extern "C" magma_int_t
magma_clarfb_gpu( magma_side_t side, magma_trans_t trans, magma_direct_t direct, magma_storev_t storev,
magma_int_t m, magma_int_t n, magma_int_t k,
const magmaFloatComplex *dV, magma_int_t ldv,
const magmaFloatComplex *dT, magma_int_t ldt,
magmaFloatComplex *dC, magma_int_t ldc,
magmaFloatComplex *dwork, magma_int_t ldwork )
{
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
/* Check input arguments */
magma_int_t info = 0;
if (m < 0) {
info = -5;
} else if (n < 0) {
info = -6;
} else if (k < 0) {
info = -7;
} else if ( ((storev == MagmaColumnwise) && (side == MagmaLeft) && ldv < max(1,m)) ||
((storev == MagmaColumnwise) && (side == MagmaRight) && ldv < max(1,n)) ||
((storev == MagmaRowwise) && ldv < k) ) {
info = -9;
} else if (ldt < k) {
info = -11;
} else if (ldc < max(1,m)) {
info = -13;
} else if ( ((side == MagmaLeft) && ldwork < max(1,n)) ||
((side == MagmaRight) && ldwork < max(1,m)) ) {
info = -15;
}
if (info != 0) {
magma_xerbla( __func__, -(info) );
return info;
}
/* Function Body */
if (m <= 0 || n <= 0) {
return info;
}
// opposite of trans
magma_trans_t transt;
if (trans == MagmaNoTrans)
transt = Magma_ConjTrans;
else
transt = MagmaNoTrans;
// whether T is upper or lower triangular
magma_uplo_t uplo;
if (direct == MagmaForward)
uplo = MagmaUpper;
else
uplo = MagmaLower;
// whether V is stored transposed or not
magma_trans_t notransV, transV;
if (storev == MagmaColumnwise) {
notransV = MagmaNoTrans;
transV = Magma_ConjTrans;
}
else {
notransV = Magma_ConjTrans;
transV = MagmaNoTrans;
}
if ( side == MagmaLeft ) {
// Form H C or H^H C
// Comments assume H C. When forming H^H C, T gets transposed via transt.
// W = C^H V
magma_cgemm( Magma_ConjTrans, notransV,
n, k, m,
c_one, dC, ldc,
dV, ldv,
c_zero, dwork, ldwork);
// W = W T^H = C^H V T^H
magma_ctrmm( MagmaRight, uplo, transt, MagmaNonUnit,
n, k,
c_one, dT, ldt,
dwork, ldwork);
// C = C - V W^H = C - V T V^H C = (I - V T V^H) C = H C
magma_cgemm( notransV, Magma_ConjTrans,
m, n, k,
c_neg_one, dV, ldv,
dwork, ldwork,
c_one, dC, ldc);
}
else {
// Form C H or C H^H
// Comments assume C H. When forming C H^H, T gets transposed via trans.
// W = C V
magma_cgemm( MagmaNoTrans, notransV,
m, k, n,
c_one, dC, ldc,
dV, ldv,
c_zero, dwork, ldwork);
// W = W T = C V T
magma_ctrmm( MagmaRight, uplo, trans, MagmaNonUnit,
m, k,
c_one, dT, ldt,
dwork, ldwork);
// C = C - W V^H = C - C V T V^H = C (I - V T V^H) = C H
magma_cgemm( MagmaNoTrans, transV,
m, n, k,
c_neg_one, dwork, ldwork,
dV, ldv,
c_one, dC, ldc);
}
return info;
} /* magma_clarfb */
/**
Purpose
-------
CGETRI computes the inverse of a matrix using the LU factorization
computed by CGETRF. 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 CGESV,
or CGETRF and CGETRS, 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 array on the GPU, dimension (LDDA,N)
On entry, the factors L and U from the factorization
A = P*L*U as computed by CGETRF_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 CGETRF; for 1 <= i <= N, row i of the
matrix was interchanged with row IPIV(i).
@param[out]
dwork (workspace) COMPLEX 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_cgetri_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_cgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_cgetri_gpu( magma_int_t n, magmaFloatComplex *dA, magma_int_t ldda,
magma_int_t *ipiv, magmaFloatComplex *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 */
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *dL = dwork;
magma_int_t lddl = n;
magma_int_t nb = magma_get_cgetri_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_ctrtri_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_clacpy( MagmaFull, n-j, jb,
dA(j,j), ldda,
dL(j,0), lddl );
magmablas_claset( 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_cgemm( 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_ctrsm( 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_cswap( n, dA(0,j), 1, dA(0,jp), 1 );
}
}
return *info;
}
/**
Purpose
-------
CLAHRU is an auxiliary MAGMA routine that is used in CGEHRD to update
the trailing sub-matrices after the reductions of the corresponding
panels.
See further details below.
Arguments
---------
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in]
ihi INTEGER
Last row to update. Same as IHI in cgehrd.
@param[in]
k INTEGER
Number of rows of the matrix Am (see details below)
@param[in]
nb INTEGER
Block size
@param[out]
A COMPLEX array, dimension (LDA,N-K)
On entry, the N-by-(N-K) general matrix to be updated. The
computation is done on the GPU. After Am is updated on the GPU
only Am(1:NB) is transferred to the CPU - to update the
corresponding Am matrix. See Further Details below.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[in,out]
data Structure with pointers to dA, dT, dV, dW, dY
which are distributed across multiple GPUs.
Further Details
---------------
This implementation follows the algorithm and notations described in:
S. Tomov and J. Dongarra, "Accelerating the reduction to upper Hessenberg
form through hybrid GPU-based computing," University of Tennessee Computer
Science Technical Report, UT-CS-09-642 (also LAPACK Working Note 219),
May 24, 2009.
The difference is that here Am is computed on the GPU.
M is renamed Am, G is renamed Ag.
@ingroup magma_cgeev_aux
********************************************************************/
extern "C" magma_int_t
magma_clahru_m(
magma_int_t n, magma_int_t ihi, magma_int_t k, magma_int_t nb,
magmaFloatComplex *A, magma_int_t lda,
struct cgehrd_data* data )
{
#define dA( d, i, j ) (data->A [d] + (i) + (j)*ldda)
#define dTi( d ) (data->Ti[d])
#define dV( d, i, j ) (data->V [d] + (i) + (j)*ldv )
#define dVd( d, i, j ) (data->Vd[d] + (i) + (j)*ldvd)
#define dW( d, i, j ) (data->W [d] + (i) + (j)*ldda)
#define dY( d, i, j ) (data->Y [d] + (i) + (j)*ldda)
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t ngpu = data->ngpu;
magma_int_t ldda = data->ldda;
magma_int_t ldv = data->ldv;
magma_int_t ldvd = data->ldvd;
magma_int_t d;
magma_int_t dk, dkhi, dknb, dn;
magma_int_t info = 0;
if (n < 0) {
info = -1;
} else if (ihi < 0 || ihi > n) {
info = -2;
} else if (k < 0 || k > n) {
info = -3;
} else if (nb < 1 || nb > n) {
info = -4;
} else if (lda < max(1,n)) {
info = -6;
}
if (info != 0) {
magma_xerbla( __func__, -(info) );
return info;
}
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
for( d = 0; d < ngpu; ++d ) {
magma_setdevice( d );
magmablasSetKernelStream( data->streams[d] );
// convert global indices (k) to local indices (dk)
magma_indices_1D_bcyclic( nb, ngpu, d, k, ihi, &dk, &dkhi );
magma_indices_1D_bcyclic( nb, ngpu, d, k+nb, n, &dknb, &dn );
// -----
// on right, A := A Q = A - A V T V'
// Update Am = Am - Am V T Vd' = Am - Ym Wd', with Wd = Vd T'
// Wd = Vd T' = V(k:ihi-1, 0:nb-1) * T(0:nb-1, 0:nb-1)'
// Vd and Wd are the portions corresponding to the block cyclic dkstribution
magma_cgemm( MagmaNoTrans, MagmaConjTrans, dkhi-dk, nb, nb,
c_one, dVd(d, dk, 0), ldvd,
dTi(d), nb,
c_zero, dW (d, dk, 0), ldda );
// Am = Am - Ym Wd' = A(0:k-1, k:ihi-1) - Ym(0:k-1, 0:nb-1) * W(k:ihi-1, 0:nb-1)'
magma_cgemm( MagmaNoTrans, MagmaConjTrans, k, dkhi-dk, nb,
c_neg_one, dY(d, 0, 0), ldda,
dW(d, dk, 0), ldda,
c_one, dA(d, 0, dk), ldda );
// -----
// on right, A := A Q = A - A V T V'
// Update Ag = Ag - Ag V T V' = Ag - Yg Wd'
// Ag = Ag - Yg Wd' = A(k:ihi-1, nb:ihi-k-1) - Y(k:ihi-1, 0:nb-1) * W(k+nb:ihi-1, 0:nb-1)'
magma_cgemm( MagmaNoTrans, MagmaConjTrans, ihi-k, dkhi-dknb, nb,
c_neg_one, dY(d, k, 0), ldda,
dW(d, dknb, 0), ldda,
c_one, dA(d, k, dknb), ldda );
// -----
// on left, A := Q' A = A - V T' V' A
// Ag2 = Ag2 - V T' V' Ag2 = W Yg, with W = V T' and Yg = V' Ag2
// Note that Ag is A(k:ihi, nb+1:ihi-k)
// while Ag2 is A(k:ihi, nb+1: n -k)
// here V and W are the whole matrices, not just block cyclic portion
// W = V T' = V(k:ihi-1, 0:nb-1) * T(0:nb-1, 0:nb-1)'
// TODO would it be cheaper to compute the whole matrix and
// copy the block cyclic portions to another workspace?
magma_cgemm( MagmaNoTrans, MagmaConjTrans, ihi-k, nb, nb,
c_one, dV (d, k, 0), ldv,
dTi(d), nb,
c_zero, dW (d, k, 0), ldda );
// Z = V(k:ihi-1, 0:nb-1)' * A(k:ihi-1, nb:n-k-1); Z is stored over Y
magma_cgemm( MagmaConjTrans, MagmaNoTrans, nb, dn-dknb, ihi-k,
c_one, dV(d, k, 0), ldv,
dA(d, k, dknb), ldda,
c_zero, dY(d, 0, 0), nb );
// Ag2 = Ag2 - W Z = A(k:ihi-1, k+nb:n-1) - W(k+nb:n-1, 0:nb-1) * Z(0:nb-1, k+nb:n-1)
magma_cgemm( MagmaNoTrans, MagmaNoTrans, ihi-k, dn-dknb, nb,
c_neg_one, dW(d, k, 0), ldda,
dY(d, 0, 0), nb,
c_one, dA(d, k, dknb), ldda );
}
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
return info;
}
/**
Purpose
-------
CGETRF_INCPIV computes an LU factorization of a general M-by-N tile A
using partial pivoting with row interchanges.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 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,out]
hA COMPLEX 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 array, dimension(LDDA, N), on gpu.
On entry, the M-by-N tile to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
@param[out]
hL COMPLEX array, dimension(LDHL, min(M,N)), on vpu.
On exit, contains in the upper part the IB-by-K lower triangular tile,
and in the lower part IB-by-min(M,N) 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 array, dimension(LDDL, K), on gpu.
On exit, contains in the upper part the IB-by-min(M,N) lower triangular tile,
and in the lower part IB-by-min(M,N) the inverse of the top part.
@param[in]
lddl INTEGER
The leading dimension of the array dL. LDDL >= max(1,2*IB).
@param[out]
ipiv INTEGER array, dimension min(M,N), on the cpu.
The pivot indices array.
@param[out]
dWORK COMPLEX 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]
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_cgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_cgetrf_incpiv_gpu( magma_order_t order, magma_int_t m, magma_int_t n, magma_int_t ib,
magmaFloatComplex *hA, magma_int_t ldha, magmaFloatComplex *dA, magma_int_t ldda,
magmaFloatComplex *hL, magma_int_t ldhl, magmaFloatComplex *dL, magma_int_t lddl,
magma_int_t *ipiv,
magmaFloatComplex *dwork, magma_int_t lddwork,
magma_int_t *info)
{
#define AT(i,j) (dAT + (i)*ib*ldda + (j)*ib)
#define hA(i,j) (hA + (i)*ib + (j)*ib*ldha)
#define hL(j) (hL + (j)*ib*ldhl )
#define hL2(j) (hL2 + (j)*ib*ldhl )
#define dL(j) (dL + (j)*ib*lddl )
#define dL2(j) (dL2 + (j)*ib*lddl )
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t iinfo;
magma_int_t maxm, mindim;
magma_int_t i, rows, cols, s, ii, sb;
magmaFloatComplex *dAT;
#ifndef WITHOUTTRTRI
magmaFloatComplex *dL2 = dL + ib;
magmaFloatComplex *hL2 = hL + ib;
#endif
/* Check arguments */
*info = 0;
if (m < 0)
*info = -1;
else if (n < 0)
*info = -2;
else if (ldda < max(1,m))
*info = -4;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return if possible */
if (m == 0 || n == 0)
return *info;
/* Function Body */
mindim = min(m, n);
s = mindim / ib;
if ( ib >= mindim ) {
/* Use CPU code. */
lapackf77_cgetrf(&m, &n, hA, &ldha, ipiv, info);
#ifndef WITHOUTTRTRI
CORE_clacpy(PlasmaUpperLower, mindim, mindim,
(PLASMA_Complex32_t*)hA, ldha,
(PLASMA_Complex32_t*)hL2, ldhl );
CORE_ctrtri( PlasmaLower, PlasmaUnit, mindim,
(PLASMA_Complex32_t*)hL2, ldhl, info );
if (*info != 0 ) {
fprintf(stderr, "ERROR, trtri returned with info = %d\n", *info);
}
magma_csetmatrix( mindim, mindim, hL2, ldhl, dL2, lddl );
#endif
if ( order == MagmaRowMajor ) {
magma_csetmatrix( m, n, hA, ldha, dwork, lddwork );
magmablas_ctranspose( m, n, dwork, lddwork, dA, ldda );
} else {
magma_csetmatrix( m, n, hA, ldha, dA, ldda );
}
}
else {
/* Use hybrid blocked code. */
maxm = ((m + 31)/32)*32;
if ( order == MagmaColMajor ) {
magmablas_cgetmo_in( dA, dAT, ldda, m, n );
} else {
dAT = dA;
}
for( i=0; i < s; i++ ) {
ii = i * ib;
sb = min(ib, mindim-ii);
cols = maxm - ii;
if ( i > 0 ) {
// download i-th panel
magmablas_ctranspose( sb, m, AT(0,i), ldda, dwork, maxm );
magma_cgetmatrix( m, sb, dwork, maxm, hA(0, i), ldha );
// make sure that gpu queue is empty
//magma_device_sync();
#ifndef WITHOUTTRTRI
magma_ctrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n - (ii+sb), ib,
c_one, dL2(i-1), lddl,
AT(i-1,i+1), ldda );
#else
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n - (ii+sb), ib,
c_one, AT(i-1,i-1), ldda,
AT(i-1,i+1), ldda );
#endif
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-(ii+sb), m-ii, ib,
c_neg_one, AT(i-1,i+1), ldda,
AT(i, i-1), ldda,
c_one, AT(i, i+1), ldda );
}
// do the cpu part
rows = m - ii;
lapackf77_cgetrf( &rows, &sb, hA(i, i), &ldha, ipiv+ii, &iinfo);
if ( (*info == 0) && (iinfo > 0) )
*info = iinfo + ii;
{
int j;
int fin = ii + sb;
for (j=ii; j < fin; j++) {
ipiv[j] = ii + ipiv[j];
}
}
magmablas_claswp( n-ii, AT(0, i), ldda, ii+1, ii+sb, ipiv, 1 );
#ifndef WITHOUTTRTRI
CORE_clacpy(PlasmaLower, sb, sb,
(PLASMA_Complex32_t*)hA(i, i), ldha,
(PLASMA_Complex32_t*)hL2(i), ldhl );
CORE_ctrtri( PlasmaLower, PlasmaUnit, sb,
(PLASMA_Complex32_t*)hL2(i), ldhl, info );
if (*info != 0 ) {
fprintf(stderr, "ERROR, trtri returned with info = %d\n", *info);
}
magma_csetmatrix( sb, sb, hL2(i), ldhl, dL2(i), lddl );
#endif
// upload i-th panel
magma_csetmatrix( rows, sb, hA(i, i), ldha, dwork, cols );
magmablas_ctranspose( rows, sb, dwork, cols, AT(i,i), ldda );
// do the small non-parallel computations
if ( s > (i+1) ) {
#ifndef WITHOUTTRTRI
magma_ctrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
sb, sb,
c_one, dL2(i), lddl,
AT(i, i+1), ldda);
#else
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
sb, sb,
c_one, AT(i, i ), ldda,
AT(i, i+1), ldda);
#endif
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
sb, m-(ii+sb), sb,
c_neg_one, AT(i, i+1), ldda,
AT(i+1, i ), ldda,
c_one, AT(i+1, i+1), ldda );
}
else {
/* Update of the last panel */
#ifndef WITHOUTTRTRI
magma_ctrmm( MagmaRight, MagmaLower, MagmaTrans, MagmaUnit,
n-mindim, sb,
c_one, dL2(i), lddl,
AT(i, i+1), ldda);
#else
magma_ctrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaUnit,
n-mindim, sb,
c_one, AT(i, i ), ldda,
AT(i, i+1), ldda);
#endif
/* m-(ii+sb) should be always 0 */
magma_cgemm( MagmaNoTrans, MagmaNoTrans,
n-mindim, m-(ii+sb), sb,
c_neg_one, AT(i, i+1), ldda,
AT(i+1, i ), ldda,
c_one, AT(i+1, i+1), ldda );
}
}
if ( order == MagmaColMajor ) {
magmablas_cgetmo_out( dA, dAT, ldda, m, n );
}
}
return *info;
}
