/*------------------------------------------------------------
* Check the reduction
*/
static magma_int_t check_reduction(magma_int_t uplo, magma_int_t N, magma_int_t bw, magmaFloatComplex *A, float *D, magma_int_t LDA, magmaFloatComplex *Q, float eps )
{
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *TEMP = (magmaFloatComplex *)malloc(N*N*sizeof(magmaFloatComplex));
magmaFloatComplex *Residual = (magmaFloatComplex *)malloc(N*N*sizeof(magmaFloatComplex));
float *work = (float *)malloc(N*sizeof(float));
float Anorm, Rnorm, result;
magma_int_t info_reduction;
magma_int_t i;
magma_int_t ione=1;
char luplo = uplo == MagmaLower ? 'L' : 'U';
/* Compute TEMP = Q * LAMBDA */
lapackf77_clacpy("A", &N, &N, Q, &LDA, TEMP, &N);
for (i = 0; i < N; i++){
blasf77_csscal(&N, &D[i], &(TEMP[i*N]), &ione);
}
/* Compute Residual = A - Q * LAMBDA * Q^H */
/* A is Hermetian but both upper and lower
* are assumed valable here for checking
* otherwise it need to be symetrized before
* checking.
*/
lapackf77_clacpy("A", &N, &N, A, &LDA, Residual, &N);
blasf77_cgemm("N", "C", &N, &N, &N, &c_neg_one, TEMP, &N, Q, &LDA, &c_one, Residual, &N);
// since A has been generated by larnv and we did not symmetrize,
// so only the uplo portion of A should be equal to Q*LAMBDA*Q^H
// for that Rnorm use clanhe instead of clange
Rnorm = lapackf77_clanhe("1", &luplo, &N, Residual, &N, work);
Anorm = lapackf77_clanhe("1", &luplo, &N, A, &LDA, work);
result = Rnorm / ( Anorm * N * eps);
if ( uplo == MagmaLower ){
printf(" ======================================================\n");
printf(" ||A-Q*LAMBDA*Q'||_oo/(||A||_oo.N.eps) : %15.3E \n", result );
printf(" ======================================================\n");
}else{
printf(" ======================================================\n");
printf(" ||A-Q'*LAMBDA*Q||_oo/(||A||_oo.N.eps) : %15.3E \n", result );
printf(" ======================================================\n");
}
if ( isnan(result) || isinf(result) || (result > 60.0) ) {
printf("-- Reduction is suspicious ! \n");
info_reduction = 1;
}
else {
printf("-- Reduction is CORRECT ! \n");
info_reduction = 0;
}
free(TEMP); free(Residual);
free(work);
return info_reduction;
}
// On input, LU and ipiv is LU factorization of A. On output, LU is overwritten.
// Works for any m, n.
// Uses init_matrix() to re-generate original A as needed.
// Returns error in factorization, |PA - LU| / (n |A|)
// This allocates 3 more matrices to store A, L, and U.
float get_LU_error(
magma_opts &opts,
magma_int_t M, magma_int_t N,
magmaFloatComplex *LU, magma_int_t lda,
magma_int_t *ipiv)
{
magma_int_t min_mn = min(M,N);
magma_int_t ione = 1;
magma_int_t i, j;
magmaFloatComplex alpha = MAGMA_C_ONE;
magmaFloatComplex beta = MAGMA_C_ZERO;
magmaFloatComplex *A, *L, *U;
float work[1], matnorm, residual;
TESTING_MALLOC_CPU( A, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( L, magmaFloatComplex, M*min_mn );
TESTING_MALLOC_CPU( U, magmaFloatComplex, min_mn*N );
memset( L, 0, M*min_mn*sizeof(magmaFloatComplex) );
memset( U, 0, min_mn*N*sizeof(magmaFloatComplex) );
// set to original A
init_matrix( opts, M, N, A, lda );
lapackf77_claswp( &N, A, &lda, &ione, &min_mn, ipiv, &ione);
// copy LU to L and U, and set diagonal to 1
lapackf77_clacpy( MagmaLowerStr, &M, &min_mn, LU, &lda, L, &M );
lapackf77_clacpy( MagmaUpperStr, &min_mn, &N, LU, &lda, U, &min_mn );
for (j=0; j < min_mn; j++)
L[j+j*M] = MAGMA_C_MAKE( 1., 0. );
matnorm = lapackf77_clange("f", &M, &N, A, &lda, work);
blasf77_cgemm("N", "N", &M, &N, &min_mn,
&alpha, L, &M, U, &min_mn, &beta, LU, &lda);
for( j = 0; j < N; j++ ) {
for( i = 0; i < M; i++ ) {
LU[i+j*lda] = MAGMA_C_SUB( LU[i+j*lda], A[i+j*lda] );
}
}
residual = lapackf77_clange("f", &M, &N, LU, &lda, work);
TESTING_FREE_CPU( A );
TESTING_FREE_CPU( L );
TESTING_FREE_CPU( U );
return residual / (matnorm * N);
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgesv
*/
int main(int argc, char **argv)
{
real_Double_t gflops, cpu_perf, cpu_time, gpu_perf, gpu_time;
float error, Rnorm, Anorm, Xnorm, *work;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_LU, *h_B, *h_X;
magma_int_t *ipiv;
magma_int_t N, nrhs, lda, ldb, info, sizeA, sizeB;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
/* Initialize */
magma_queue_t queue[2];
magma_device_t device[ MagmaMaxGPUs ];
int num = 0;
magma_err_t err;
magma_init();
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
nrhs = opts.nrhs;
err = magma_get_devices( device, MagmaMaxGPUs, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_get_devices failed: %d\n", err );
exit(-1);
}
// Create two queues on device opts.device
err = magma_queue_create( device[ opts.device ], &queue[0] );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( device[ opts.device ], &queue[1] );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
printf("ngpu %d\n", (int) opts.ngpu );
printf(" N NRHS CPU Gflop/s (sec) GPU GFlop/s (sec) ||B - AX|| / N*||A||*||X||\n");
printf("================================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
lda = N;
ldb = lda;
gflops = ( FLOPS_CGETRF( N, N ) + FLOPS_CGETRS( N, nrhs ) ) / 1e9;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( h_LU, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, ldb*nrhs );
TESTING_MALLOC_CPU( h_X, magmaFloatComplex, ldb*nrhs );
TESTING_MALLOC_CPU( work, float, N );
TESTING_MALLOC_CPU( ipiv, magma_int_t, N );
/* Initialize the matrices */
sizeA = lda*N;
sizeB = ldb*nrhs;
lapackf77_clarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_clarnv( &ione, ISEED, &sizeB, h_B );
// copy A to LU and B to X; save A and B for residual
lapackf77_clacpy( "F", &N, &N, h_A, &lda, h_LU, &lda );
lapackf77_clacpy( "F", &N, &nrhs, h_B, &ldb, h_X, &ldb );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_cgesv( N, nrhs, h_LU, lda, ipiv, h_X, ldb, &info, queue );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgesv returned error %d: %s.\n",
(int) info, magma_strerror( info ));
//=====================================================================
// Residual
//=====================================================================
Anorm = lapackf77_clange("I", &N, &N, h_A, &lda, work);
Xnorm = lapackf77_clange("I", &N, &nrhs, h_X, &ldb, work);
blasf77_cgemm( MagmaNoTransStr, MagmaNoTransStr, &N, &nrhs, &N,
&c_one, h_A, &lda,
h_X, &ldb,
&c_neg_one, h_B, &ldb);
Rnorm = lapackf77_clange("I", &N, &nrhs, h_B, &ldb, work);
error = Rnorm/(N*Anorm*Xnorm);
status |= ! (error < tol);
/* ====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_cgesv( &N, &nrhs, h_A, &lda, ipiv, h_B, &ldb, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cgesv returned error %d: %s.\n",
(int) info, magma_strerror( info ));
printf( "%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e%s\n",
(int) N, (int) nrhs, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "" : " failed"));
}
else {
printf( "%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e%s\n",
(int) N, (int) nrhs, gpu_perf, gpu_time,
error, (error < tol ? "" : " failed"));
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_LU );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_X );
TESTING_FREE_CPU( work );
TESTING_FREE_CPU( ipiv );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
magma_queue_destroy( queue[0] );
magma_queue_destroy( queue[1] );
magma_finalize();
return status;
}
int main( int argc, char** argv )
{
magma_init();
cublasHandle_t handle;
cudaSetDevice( 0 );
cublasCreate( &handle );
magmaFloatComplex *A, *B, *C;
magmaFloatComplex *dA, *dB, *dC;
float error, work[1];
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = { 1, 2, 3, 4 };
magma_int_t n = 10;
magma_int_t lda = n;
magma_int_t ldda = ((n+31)/32)*32;
magma_int_t size = lda*n;
magma_int_t info;
magma_cmalloc_cpu( &A, lda*n );
magma_cmalloc_cpu( &B, lda*n );
magma_cmalloc_cpu( &C, lda*n );
magma_cmalloc( &dA, ldda*n );
magma_cmalloc( &dB, ldda*n );
magma_cmalloc( &dC, ldda*n );
// initialize matrices
lapackf77_clarnv( &ione, ISEED, &size, A );
lapackf77_clarnv( &ione, ISEED, &size, B );
lapackf77_clarnv( &ione, ISEED, &size, C );
// increase diagonal to be SPD
for( int i=0; i < n; ++i ) {
C[i+i*lda] = MAGMA_C_ADD( C[i+i*lda], MAGMA_C_MAKE( n*n, 0 ));
}
magma_csetmatrix( n, n, A, lda, dA, ldda );
magma_csetmatrix( n, n, B, lda, dB, ldda );
magma_csetmatrix( n, n, C, lda, dC, ldda );
// compute with cublas
cublasCgemm( handle, CUBLAS_OP_N, CUBLAS_OP_N, n, n, n,
&c_neg_one, dA, ldda, dB, ldda, &c_one, dC, ldda );
magma_cpotrf_gpu( MagmaLower, n, dC, ldda, &info );
if (info != 0)
printf("magma_cpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// compute with LAPACK
blasf77_cgemm( MagmaNoTransStr, MagmaNoTransStr, &n, &n, &n,
&c_neg_one, A, &lda, B, &lda, &c_one, C, &lda );
lapackf77_cpotrf( MagmaLowerStr, &n, C, &lda, &info );
if (info != 0)
printf("lapackf77_cpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// compute difference
magma_cgetmatrix( n, n, dC, ldda, A, lda );
blasf77_caxpy( &size, &c_neg_one, C, &ione, A, &ione );
error = lapackf77_clange( "F", &n, &n, A, &lda, work );
printf( "n %d, error %8.2e\n", (int) n, error );
magma_free( dA );
magma_free( dB );
magma_free( dC );
magma_free_cpu( A );
magma_free_cpu( B );
magma_free_cpu( C );
cublasDestroy( handle );
magma_finalize();
return 0;
}
/***************************************************************************//**
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]
dA COMPLEX array, dimension (LDA,N)
Copy of A on the GPU.
Portions of A are updated on the CPU;
portions of dA are updated on the GPU. See code for details.
@param[in]
ldda INTEGER
The leading dimension of the array dA. 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]
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).
@param[in,out]
dF COMPLEX array, dimension (LDDF,NB)
Copy of F on the GPU. See code for details.
@param[in]
lddf INTEGER
The leading dimension of the array dF. LDDF >= max(1,N).
@ingroup magma_laqps
*******************************************************************************/
extern "C" magma_int_t
magma_claqps(
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,
magmaFloatComplex_ptr dA, magma_int_t ldda,
magma_int_t *jpvt, magmaFloatComplex *tau, float *vn1, float *vn2,
magmaFloatComplex *auxv,
magmaFloatComplex *F, magma_int_t ldf,
magmaFloatComplex_ptr dF, magma_int_t lddf)
{
#define A(i, j) (A + (i) + (j)*(lda ))
#define dA(i, j) (dA + (i) + (j)*(ldda))
#define F(i, j) (F + (i) + (j)*(ldf ))
#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;
float d__1;
magmaFloatComplex z__1;
magma_int_t j, k, rk;
magmaFloatComplex Akk;
magma_int_t pvt;
float temp, temp2, tol3z;
magma_int_t itemp;
magma_int_t lsticc;
magma_int_t lastrk;
lastrk = min( m, n + offset );
tol3z = magma_ssqrt( lapackf77_slamch("Epsilon"));
magma_queue_t queue;
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queue );
lsticc = 0;
k = 0;
while( k < nb && lsticc == 0 ) {
rk = offset + k;
/* Determine ith pivot column and swap if necessary */
// subtract 1 from Fortran isamax; pvt, k are 0-based.
i__1 = n-k;
pvt = k + blasf77_isamax( &i__1, &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, queue );
}
/* F gets swapped so F must be sent at the end to GPU */
i__1 = k;
blasf77_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 (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( queue );
/* 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, 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) {
#ifdef COMPLEX
for (j = 0; j < k; ++j) {
*F(k,j) = MAGMA_C_CONJ( *F(k,j) );
}
#endif
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 );
#ifdef COMPLEX
for (j = 0; j < k; ++j) {
*F(k,j) = MAGMA_C_CONJ( *F(k,j) );
}
#endif
}
/* Generate elementary reflector H(k). */
if (rk < m-1) {
i__1 = m - rk;
lapackf77_clarfg( &i__1, A(rk, k), A(rk + 1, k), &ione, &tau[k] );
} else {
lapackf77_clarfg( &ione, A(rk, k), A(rk, k), &ione, &tau[k] );
}
Akk = *A(rk, k);
*A(rk, k) = c_one;
/* 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) {
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, queue );
/* 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_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, queue );
magma_cgetmatrix_async( i__2-i__3, 1,
dF(k + 1 +i__3, k), i__2,
F (k + 1 +i__3, k), i__2, queue );
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( queue );
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) {
*F(j, k) = c_zero;
}
/* 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). */
if (k > 0) {
i__1 = m - rk;
i__2 = k;
z__1 = MAGMA_C_NEGATE( tau[k] );
blasf77_cgemv( MagmaConjTransStr, &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 );
}
/* 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 );
}
/* Update partial column norms. */
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;
++k;
}
// 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, queue );
magma_cgemm( MagmaNoTrans, MagmaConjTrans, i__1, i__2, *kb,
c_neg_one, dA(rk+1, 0 ), ldda,
dF(*kb, 0 ), i__2,
c_one, dA(rk+1, *kb), ldda, queue );
}
/* Recomputation of difficult columns. */
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, queue );
//vn1[lsticc] = magma_scnrm2( i__1, dA(rk + 1, lsticc), ione, queue );
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_queue_destroy( queue );
return MAGMA_SUCCESS;
} /* magma_claqps */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgemm_batched
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cublas_perf, cublas_time, cpu_perf, cpu_time;
float magma_error, cublas_error, magma_err, cublas_err, 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;
magma_int_t NN;
magma_int_t batchCount;
magmaFloatComplex *h_A, *h_B, *h_C, *h_Cmagma, *h_Ccublas;
magmaFloatComplex *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 );
magmaFloatComplex **A_array = NULL;
magmaFloatComplex **B_array = NULL;
magmaFloatComplex **C_array = NULL;
magma_queue_t queue = magma_stream;
magma_opts opts;
parse_opts( argc, argv, &opts );
batchCount = opts.batchcount;
cublasHandle_t handle = opts.handle;
//float tol = opts.tolerance * lapackf77_slamch("E");
printf("If running lapack (option --lapack), MAGMA and CUBLAS error are both computed\n"
"relative to CPU BLAS result. Else, MAGMA error is computed relative to CUBLAS result.\n\n"
"transA = %s, transB = %s\n",
lapack_trans_const(opts.transA),
lapack_trans_const(opts.transB));
printf("BatchCount M N K MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) CPU Gflop/s (ms) MAGMA error CUBLAS error\n");
printf("=========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
K = opts.ksize[itest];
gflops = FLOPS_CGEMM( M, N, K ) / 1e9 * batchCount;
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;
NN = N * batchCount;
ldda = ((lda+31)/32)*32;
lddb = ((ldb+31)/32)*32;
lddc = ((ldc+31)/32)*32;
sizeA = lda*An*batchCount;
sizeB = ldb*Bn*batchCount;
sizeC = ldc*N*batchCount;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, sizeA );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, sizeB );
TESTING_MALLOC_CPU( h_C, magmaFloatComplex, sizeC );
TESTING_MALLOC_CPU( h_Cmagma, magmaFloatComplex, sizeC );
TESTING_MALLOC_CPU( h_Ccublas, magmaFloatComplex, sizeC );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*An*batchCount );
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, lddb*Bn*batchCount );
TESTING_MALLOC_DEV( d_C, magmaFloatComplex, lddc*N*batchCount );
magma_malloc((void**)&A_array, batchCount * sizeof(*A_array));
magma_malloc((void**)&B_array, batchCount * sizeof(*B_array));
magma_malloc((void**)&C_array, batchCount * sizeof(*C_array));
/* Initialize the matrices */
lapackf77_clarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_clarnv( &ione, ISEED, &sizeB, h_B );
lapackf77_clarnv( &ione, ISEED, &sizeC, h_C );
/* =====================================================================
Performs operation using MAGMABLAS
=================================================================== */
magma_csetmatrix( Am, An*batchCount, h_A, lda, d_A, ldda );
magma_csetmatrix( Bm, Bn*batchCount, h_B, ldb, d_B, lddb );
magma_csetmatrix( M, N*batchCount, h_C, ldc, d_C, lddc );
cset_pointer(A_array, d_A, ldda, 0, 0, ldda*An, batchCount, queue);
cset_pointer(B_array, d_B, lddb, 0, 0, lddb*Bn, batchCount, queue);
cset_pointer(C_array, d_C, lddc, 0, 0, lddc*N, batchCount, queue);
magma_time = magma_sync_wtime( NULL );
magmablas_cgemm_batched(opts.transA, opts.transB, M, N, K,
alpha, A_array, ldda,
B_array, lddb,
beta, C_array, lddc, batchCount, queue);
magma_time = magma_sync_wtime( NULL ) - magma_time;
magma_perf = gflops / magma_time;
magma_cgetmatrix( M, N*batchCount, d_C, lddc, h_Cmagma, ldc );
/* =====================================================================
Performs operation using CUBLAS
=================================================================== */
magma_csetmatrix( M, N*batchCount, h_C, ldc, d_C, lddc );
cublas_time = magma_sync_wtime( NULL );
cublasCgemmBatched(handle, cublas_trans_const(opts.transA), cublas_trans_const(opts.transB), M, N, K,
&alpha, (const magmaFloatComplex**) A_array, ldda,
(const magmaFloatComplex**) B_array, lddb,
&beta, C_array, lddc, batchCount );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_cgetmatrix( M, N*batchCount, d_C, lddc, h_Ccublas, ldc );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
for(int i=0; i<batchCount; i++)
{
blasf77_cgemm(
lapack_trans_const(opts.transA), lapack_trans_const(opts.transB),
&M, &N, &K,
&alpha, h_A + i*lda*An, &lda,
h_B + i*ldb*Bn, &ldb,
&beta, h_C + i*ldc*N, &ldc );
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
if ( opts.lapack ) {
// compute relative error for both magma & cublas, relative to lapack,
// |C_magma - C_lapack| / |C_lapack|
magma_error = 0.0;
cublas_error = 0.0;
for(int s=0; s<batchCount; s++)
{
magma_int_t C_batchSize = ldc * N;
Cnorm = lapackf77_clange( "M", &M, &N, h_C + s*C_batchSize, &ldc, work );
blasf77_caxpy( &C_batchSize, &c_neg_one, h_C + s*C_batchSize, &ione, h_Cmagma + s*C_batchSize, &ione );
magma_err = lapackf77_clange( "M", &M, &N, h_Cmagma + s*C_batchSize, &ldc, work ) / Cnorm;
if ( isnan(magma_err) || isinf(magma_err) ) {
magma_error = magma_err;
break;
}
magma_error = max(fabs(magma_err), magma_error);
blasf77_caxpy( &C_batchSize, &c_neg_one, h_C + s*C_batchSize, &ione, h_Ccublas + s*C_batchSize, &ione );
cublas_err = lapackf77_clange( "M", &M, &N, h_Ccublas + s*C_batchSize, &ldc, work ) / Cnorm;
if ( isnan(cublas_err) || isinf(cublas_err) ) {
cublas_error = cublas_err;
break;
}
cublas_error = max(fabs(cublas_err), cublas_error);
}
printf("%10d %5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e \n",
(int) batchCount, (int) M, (int) N, (int) K,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error);
}
else {
// compute relative error for magma, relative to cublas
Cnorm = lapackf77_clange( "M", &M, &NN, h_Ccublas, &ldc, work );
blasf77_caxpy( &sizeC, &c_neg_one, h_Ccublas, &ione, h_Cmagma, &ione );
magma_error = lapackf77_clange( "M", &M, &NN, h_Cmagma, &ldc, work ) / Cnorm;
printf("%10d %5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) --- ( --- ) %8.2e ---\n",
(int) batchCount, (int) M, (int) N, (int) K,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
magma_error );
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_C );
TESTING_FREE_CPU( h_Cmagma );
TESTING_FREE_CPU( h_Ccublas );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
TESTING_FREE_DEV( d_C );
TESTING_FREE_DEV( A_array );
TESTING_FREE_DEV( B_array );
TESTING_FREE_DEV( C_array );
fflush( stdout);
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cposv
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, cpu_perf, cpu_time, gpu_perf, gpu_time;
float error, Rnorm, Anorm, Xnorm, *work;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_R, *h_B, *h_X;
magma_int_t N, lda, ldb, info, sizeA, sizeB;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf("ngpu %d, uplo %c\n", (int) opts.ngpu, opts.uplo );
printf(" N NRHS CPU Gflop/s (sec) GPU GFlop/s (sec) ||B - AX|| / N*||A||*||X||\n");
printf("================================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[i];
lda = ldb = N;
gflops = ( FLOPS_CPOTRF( N ) + FLOPS_CPOTRS( N, opts.nrhs ) ) / 1e9;
TESTING_MALLOC( h_A, magmaFloatComplex, lda*N );
TESTING_MALLOC( h_R, magmaFloatComplex, lda*N );
TESTING_MALLOC( h_B, magmaFloatComplex, ldb*opts.nrhs );
TESTING_MALLOC( h_X, magmaFloatComplex, ldb*opts.nrhs );
TESTING_MALLOC( work, float, N );
/* ====================================================================
Initialize the matrix
=================================================================== */
sizeA = lda*N;
sizeB = ldb*opts.nrhs;
lapackf77_clarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_clarnv( &ione, ISEED, &sizeB, h_B );
magma_cmake_hpd( N, h_A, lda );
// copy A to R and B to X; save A and B for residual
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &lda, h_R, &lda );
lapackf77_clacpy( MagmaUpperLowerStr, &N, &opts.nrhs, h_B, &ldb, h_X, &ldb );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_cposv( opts.uplo, N, opts.nrhs, h_R, lda, h_X, ldb, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Residual
=================================================================== */
Anorm = lapackf77_clange("I", &N, &N, h_A, &lda, work);
Xnorm = lapackf77_clange("I", &N, &opts.nrhs, h_X, &ldb, work);
blasf77_cgemm( MagmaNoTransStr, MagmaNoTransStr, &N, &opts.nrhs, &N,
&c_one, h_A, &lda,
h_X, &ldb,
&c_neg_one, h_B, &ldb );
Rnorm = lapackf77_clange("I", &N, &opts.nrhs, h_B, &ldb, work);
error = Rnorm/(N*Anorm*Xnorm);
status |= ! (error < tol);
/* ====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_cposv( &opts.uplo, &N, &opts.nrhs, h_A, &lda, h_B, &ldb, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cposv returned error %d: %s.\n",
(int) info, magma_strerror( info ));
printf( "%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e%s\n",
(int) N, (int) opts.nrhs, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "" : " failed"));
}
else {
printf( "%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e%s\n",
(int) N, (int) opts.nrhs, gpu_perf, gpu_time,
error, (error < tol ? "" : " failed"));
}
TESTING_FREE( h_A );
TESTING_FREE( h_R );
TESTING_FREE( h_B );
TESTING_FREE( h_X );
TESTING_FREE( work );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgelqf_gpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
const float d_neg_one = MAGMA_D_NEG_ONE;
const float d_one = MAGMA_D_ONE;
const magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
const magmaFloatComplex c_one = MAGMA_C_ONE;
const magmaFloatComplex c_zero = MAGMA_C_ZERO;
const magma_int_t ione = 1;
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
float Anorm, error=0, error2=0;
magmaFloatComplex *h_A, *h_R, *tau, *h_work, tmp[1];
magmaFloatComplex_ptr d_A;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn, nb;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
float tol = opts.tolerance * lapackf77_slamch("E");
printf("%% M N CPU Gflop/s (sec) GPU Gflop/s (sec) |L - A*Q^H| |I - Q*Q^H|\n");
printf("%%==============================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
ldda = magma_roundup( M, opts.align ); // multiple of 32 by default
n2 = lda*N;
nb = magma_get_cgeqrf_nb( M, N );
gflops = FLOPS_CGELQF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_cgelqf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
lwork = max( lwork, M*nb );
TESTING_MALLOC_CPU( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, lwork );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaFullStr, &M, &N, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_csetmatrix( M, N, h_R, lda, d_A, ldda, opts.queue );
gpu_time = magma_wtime();
magma_cgelqf_gpu( M, N, d_A, ldda, tau, h_work, lwork, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_cgelqf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Check the result, following zlqt01 except using the reduced Q.
This works for any M,N (square, tall, wide).
=================================================================== */
if ( opts.check ) {
magma_cgetmatrix( M, N, d_A, ldda, h_R, lda, opts.queue );
magma_int_t ldq = min_mn;
magma_int_t ldl = M;
magmaFloatComplex *Q, *L;
float *work;
TESTING_MALLOC_CPU( Q, magmaFloatComplex, ldq*N ); // K by N
TESTING_MALLOC_CPU( L, magmaFloatComplex, ldl*min_mn ); // M by K
TESTING_MALLOC_CPU( work, float, min_mn );
// generate K by N matrix Q, where K = min(M,N)
lapackf77_clacpy( "Upper", &min_mn, &N, h_R, &lda, Q, &ldq );
lapackf77_cunglq( &min_mn, &N, &min_mn, Q, &ldq, tau, h_work, &lwork, &info );
assert( info == 0 );
// copy N by K matrix L
lapackf77_claset( "Upper", &M, &min_mn, &c_zero, &c_zero, L, &ldl );
lapackf77_clacpy( "Lower", &M, &min_mn, h_R, &lda, L, &ldl );
// error = || L - A*Q^H || / (N * ||A||)
blasf77_cgemm( "NoTrans", "Conj", &M, &min_mn, &N,
&c_neg_one, h_A, &lda, Q, &ldq, &c_one, L, &ldl );
Anorm = lapackf77_clange( "1", &M, &N, h_A, &lda, work );
error = lapackf77_clange( "1", &M, &min_mn, L, &ldl, work );
if ( N > 0 && Anorm > 0 )
error /= (N*Anorm);
// set L = I (K by K), then L = I - Q*Q^H
// error = || I - Q*Q^H || / N
lapackf77_claset( "Upper", &min_mn, &min_mn, &c_zero, &c_one, L, &ldl );
blasf77_cherk( "Upper", "NoTrans", &min_mn, &N, &d_neg_one, Q, &ldq, &d_one, L, &ldl );
error2 = safe_lapackf77_clanhe( "1", "Upper", &min_mn, L, &ldl, work );
if ( N > 0 )
error2 /= N;
TESTING_FREE_CPU( Q ); Q = NULL;
TESTING_FREE_CPU( L ); L = NULL;
TESTING_FREE_CPU( work ); work = NULL;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_cgelqf( &M, &N, h_A, &lda, tau, h_work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapack_cgelqf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
}
/* =====================================================================
Print performance and error.
=================================================================== */
printf("%5d %5d ", (int) M, (int) N );
if ( opts.lapack ) {
printf( "%7.2f (%7.2f)", cpu_perf, cpu_time );
}
else {
printf(" --- ( --- )" );
}
printf( " %7.2f (%7.2f) ", gpu_perf, gpu_time );
if ( opts.check ) {
bool okay = (error < tol && error2 < tol);
printf( "error %.4g, error2 %.4g, tol %.4g, okay %d\n", error, error2, tol, okay );
status += ! okay;
printf( "%11.2e %11.2e %s\n", error, error2, (okay ? "ok" : "failed") );
}
else {
printf( " ---\n" );
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_PIN( h_R );
TESTING_FREE_PIN( h_work );
TESTING_FREE_DEV( d_A );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgegqr
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float e1, e2, e3, e4, e5, *work;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
magmaFloatComplex *h_A, *h_R, *tau, *dtau, *h_work, *h_rwork, tmp[1];
magmaFloatComplex *d_A, *dwork;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn;
magma_int_t ione = 1, ldwork;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
// versions 1...4 are valid
if (opts.version < 1 || opts.version > 4) {
printf("Unknown version %d; exiting\n", opts.version );
return -1;
}
float tol, eps = lapackf77_slamch("E");
tol = 10* opts.tolerance * eps;
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||I-Q'Q||_F / M ||I-Q'Q||_I / M ||A-Q R||_I\n");
printf(" MAGMA / LAPACK MAGMA / LAPACK\n");
printf("==========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
if (N > 128) {
printf("%5d %5d skipping because cgegqr requires N <= 128\n",
(int) M, (int) N);
continue;
}
if (M < N) {
printf("%5d %5d skipping because cgegqr requires M >= N\n",
(int) M, (int) N);
continue;
}
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_CGEQRF( M, N ) / 1e9 + FLOPS_CUNGQR( M, N, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_cgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
lwork = max(lwork, 3*N*N);
ldwork = N*N;
if (opts.version == 2) {
ldwork = 3*N*N + min_mn;
}
TESTING_MALLOC_PIN( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, lwork );
TESTING_MALLOC_PIN(h_rwork, magmaFloatComplex, lwork );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( work, float, M );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( dtau, magmaFloatComplex, min_mn );
TESTING_MALLOC_DEV( dwork, magmaFloatComplex, ldwork );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
// warmup
magma_cgegqr_gpu( 1, M, N, d_A, ldda, dwork, h_work, &info );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime( 0 );
magma_cgegqr_gpu( opts.version, M, N, d_A, ldda, dwork, h_rwork, &info );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgegqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
magma_cgetmatrix( M, N, d_A, ldda, h_R, M );
// Regenerate R
// blasf77_cgemm("t", "n", &N, &N, &M, &c_one, h_R, &M, h_A, &M, &c_zero, h_rwork, &N);
// magma_cprint(N, N, h_work, N);
blasf77_ctrmm("r", "u", "n", "n", &M, &N, &c_one, h_rwork, &N, h_R, &M);
blasf77_caxpy( &n2, &c_neg_one, h_A, &ione, h_R, &ione );
e5 = lapackf77_clange("i", &M, &N, h_R, &M, work) /
lapackf77_clange("i", &M, &N, h_A, &lda, work);
magma_cgetmatrix( M, N, d_A, ldda, h_R, M );
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
/* Orthogonalize on the CPU */
lapackf77_cgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
lapackf77_cungqr(&M, &N, &N, h_A, &lda, tau, h_work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cungqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
blasf77_cgemm("t", "n", &N, &N, &M, &c_one, h_R, &M, h_R, &M, &c_zero, h_work, &N);
for(int ii = 0; ii < N*N; ii += N+1 ) {
h_work[ii] = MAGMA_C_SUB(h_work[ii], c_one);
}
e1 = lapackf77_clange("f", &N, &N, h_work, &N, work) / N;
e3 = lapackf77_clange("i", &N, &N, h_work, &N, work) / N;
blasf77_cgemm("t", "n", &N, &N, &M, &c_one, h_A, &M, h_A, &M, &c_zero, h_work, &N);
for(int ii = 0; ii < N*N; ii += N+1 ) {
h_work[ii] = MAGMA_C_SUB(h_work[ii], c_one);
}
e2 = lapackf77_clange("f", &N, &N, h_work, &N, work) / N;
e4 = lapackf77_clange("i", &N, &N, h_work, &N, work) / N;
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e / %8.2e %8.2e / %8.2e %8.2e %s\n",
(int) M, (int) N, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time,
e1, e2, e3, e4, e5,
(e1 < tol ? "ok" : "failed"));
status += ! (e1 < tol);
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, 1000.*gpu_time );
}
TESTING_FREE_PIN( tau );
TESTING_FREE_PIN( h_work );
TESTING_FREE_PIN( h_rwork );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_R );
TESTING_FREE_CPU( work );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dtau );
TESTING_FREE_DEV( dwork );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
CGETRF_NOPIV computes an LU factorization of a general M-by-N
matrix A without pivoting.
The factorization has the form
A = L * U
where 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]
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.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,M).
@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 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(
magma_int_t m, magma_int_t n,
magmaFloatComplex *A, magma_int_t lda,
magma_int_t *info)
{
#define A(i_,j_) (A + (i_) + (j_)*lda)
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t min_mn, i__3, i__4;
magma_int_t j, jb, nb, iinfo;
A -= 1 + lda;
/* Function Body */
*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;
}
/* Determine the block size for this environment. */
nb = 128;
min_mn = min(m,n);
if (nb <= 1 || nb >= min_mn) {
/* Use unblocked code. */
magma_cgetf2_nopiv( m, n, A(1,1), lda, info );
}
else {
/* Use blocked code. */
for (j = 1; j <= min_mn; j += nb) {
jb = min( min_mn - j + 1, nb );
/* Factor diagonal and subdiagonal blocks and test for exact
singularity. */
i__3 = m - j + 1;
//magma_cgetf2_nopiv( i__3, jb, A(j,j), lda, &iinfo );
i__3 -= jb;
magma_cgetf2_nopiv( jb, jb, A(j,j), lda, &iinfo );
blasf77_ctrsm( "R", "U", "N", "N", &i__3, &jb, &c_one,
A(j,j), &lda,
A(j+jb,j), &lda );
/* Adjust INFO */
if (*info == 0 && iinfo > 0)
*info = iinfo + j - 1;
if (j + jb <= n) {
/* Compute block row of U. */
i__3 = n - j - jb + 1;
blasf77_ctrsm( "Left", "Lower", "No transpose", "Unit",
&jb, &i__3, &c_one,
A(j,j), &lda,
A(j,j+jb), &lda );
if (j + jb <= m) {
/* Update trailing submatrix. */
i__3 = m - j - jb + 1;
i__4 = n - j - jb + 1;
blasf77_cgemm( "No transpose", "No transpose",
&i__3, &i__4, &jb, &c_neg_one,
A(j+jb,j), &lda,
A(j,j+jb), &lda, &c_one,
A(j+jb,j+jb), &lda );
}
}
}
}
return *info;
} /* magma_cgetrf_nopiv */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgegqr
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float e1, e2, work[1];
magmaFloatComplex *h_A, *h_R, *tau, *dtau, *h_work, tmp[1];
magmaFloatComplex *d_A, *dwork, *ddA, *d_T;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
opts.lapack |= opts.check; // check (-c) implies lapack (-l)
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||I - Q'Q||_F \n");
printf("=======================================================================\n");
for( int i = 0; i < opts.ntest; ++i ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[i];
N = opts.nsize[i];
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = FLOPS_CGEQRF( M, N ) / 1e9 + FLOPS_CUNGQR( M, N, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_cgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_C_REAL( tmp[0] );
lwork = max(lwork, 3*N*N);
TESTING_MALLOC_PIN( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_PIN( h_work, magmaFloatComplex, lwork );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, n2 );
TESTING_MALLOC_CPU( h_R, magmaFloatComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( dtau, magmaFloatComplex, min_mn );
TESTING_MALLOC_DEV( dwork, magmaFloatComplex, N*N );
TESTING_MALLOC_DEV( ddA, magmaFloatComplex, N*N );
TESTING_MALLOC_DEV( d_T, magmaFloatComplex, N*N );
cudaMemset( ddA, 0, N*N*sizeof(magmaFloatComplex) );
cudaMemset( d_T, 0, N*N*sizeof(magmaFloatComplex) );
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
// warmup
magma_cgegqr_gpu( M, N, d_A, ldda, dwork, h_work, &info );
magma_csetmatrix( M, N, h_R, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime( 0 );
if (opts.version == 2) {
int min_mn = min(M, N);
int nb = N;
cuFloatComplex *dtau = dwork;
magma_cgeqr2x3_gpu(&M, &N, d_A, &ldda, dtau, d_T, ddA,
(float *)(dwork+min_mn), &info);
magma_cgetmatrix( min_mn, 1, dtau, min_mn, tau, min_mn);
magma_cungqr_gpu( M, N, N, d_A, ldda, tau, d_T, nb, &info );
}
else
magma_cgegqr_gpu( M, N, d_A, ldda, dwork, h_work, &info );
gpu_time = magma_sync_wtime( 0 ) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgegqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.lapack ) {
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
/* Orthogonalize on the CPU */
lapackf77_cgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
lapackf77_cungqr(&M, &N, &N, h_A, &lda, tau, h_work, &lwork, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cungqr returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
magma_cgetmatrix( M, N, d_A, ldda, h_R, M );
magmaFloatComplex one = MAGMA_C_ONE, zero = MAGMA_C_ZERO;
blasf77_cgemm("t", "n", &N, &N, &M, &one, h_R, &M, h_R, &M, &zero, h_work, &N);
for(int ii=0; ii<N*N; ii+=(N+1)) h_work[ii] = MAGMA_C_SUB(h_work[ii], one);
e1 = lapackf77_clange("f", &N, &N, h_work, &N, work);
blasf77_cgemm("t", "n", &N, &N, &M, &one, h_A, &M, h_A, &M, &zero, h_work, &N);
for(int ii=0; ii<N*N; ii+=(N+1)) h_work[ii] = MAGMA_C_SUB(h_work[ii], one);
e2 = lapackf77_clange("f", &N, &N, h_work, &N, work);
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e\n",
(int) M, (int) N, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time, e1, e2 );
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, 1000.*gpu_time );
}
TESTING_FREE_PIN( tau );
TESTING_FREE_PIN( h_work );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dtau );
TESTING_FREE_DEV( dwork );
TESTING_FREE_DEV( ddA );
TESTING_FREE_DEV( d_T );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgesv_rbt
*/
int main(int argc, char **argv)
{
TESTING_INIT();
real_Double_t gflops, cpu_perf, cpu_time, gpu_perf, gpu_time;
float error, Rnorm, Anorm, Xnorm, *work;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_LU, *h_B, *h_X;
magma_int_t *ipiv;
magma_int_t N, nrhs, lda, ldb, info, sizeA, sizeB;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
float tol = opts.tolerance * lapackf77_slamch("E");
nrhs = opts.nrhs;
printf("%% ngpu %d\n", (int) opts.ngpu );
printf("%% N NRHS CPU Gflop/s (sec) GPU Gflop/s (sec) ||B - AX|| / N*||A||*||X||\n");
printf("%%===============================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
ldb = lda;
gflops = ( FLOPS_CGETRF( N, N ) + FLOPS_CGETRS( N, nrhs ) ) / 1e9;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( h_LU, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, ldb*nrhs );
TESTING_MALLOC_CPU( h_X, magmaFloatComplex, ldb*nrhs );
TESTING_MALLOC_CPU( work, float, N );
TESTING_MALLOC_CPU( ipiv, magma_int_t, N );
/* Initialize the matrices */
sizeA = lda*N;
sizeB = ldb*nrhs;
lapackf77_clarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_clarnv( &ione, ISEED, &sizeB, h_B );
// copy A to LU and B to X; save A and B for residual
lapackf77_clacpy( "F", &N, &N, h_A, &lda, h_LU, &lda );
lapackf77_clacpy( "F", &N, &nrhs, h_B, &ldb, h_X, &ldb );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_cgesv_rbt(MagmaTrue, N, nrhs, h_LU, lda, h_X, ldb, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_cgesv_rbt returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
for (int i = 0; i < N; i++)
ipiv[i] = i+1;
//=====================================================================
// Residual
//=====================================================================
Anorm = lapackf77_clange("I", &N, &N, h_A, &lda, work);
Xnorm = lapackf77_clange("I", &N, &nrhs, h_X, &ldb, work);
blasf77_cgemm( MagmaNoTransStr, MagmaNoTransStr, &N, &nrhs, &N,
&c_one, h_A, &lda,
h_X, &ldb,
&c_neg_one, h_B, &ldb);
Rnorm = lapackf77_clange("I", &N, &nrhs, h_B, &ldb, work);
error = Rnorm/(N*Anorm*Xnorm);
status += ! (error < tol);
/* ====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_cgesv( &N, &nrhs, h_A, &lda, ipiv, h_B, &ldb, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_cgesv returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
printf( "%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) N, (int) nrhs, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed"));
}
else {
printf( "%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e %s\n",
(int) N, (int) nrhs, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed"));
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_LU );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_X );
TESTING_FREE_CPU( work );
TESTING_FREE_CPU( ipiv );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/**
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 = MagmaUpper, or
dA = L * L**H, if UPLO = MagmaLower,
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of dA is stored;
- = MagmaLower: Lower triangle of dA is stored.
@param[in]
n INTEGER
The order of the matrix dA. N >= 0.
@param[in,out]
d_lA COMPLEX array of pointers on the GPU, dimension (ngpu)
On entry, the Hermitian matrix dA distributed over GPUs
(dl_A[d] points to the local matrix on the d-th GPU).
It is distributed in 1D block column or row cyclic (with the
block size of nb) if UPLO = MagmaUpper or MagmaLower, respectively.
If UPLO = MagmaUpper, the leading N-by-N upper triangular
part of dA contains the upper triangular part of the matrix dA,
and the strictly lower triangular part of dA is not referenced.
If UPLO = MagmaLower, the leading N-by-N lower triangular part
of dA contains the lower triangular part of the matrix dA, and
the strictly upper triangular part of dA is not referenced.
\n
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization dA = U**H * U or dA = L * L**H.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
@ingroup magma_cposv_comp
********************************************************************/
extern "C" magma_int_t
magma_cpotrf_mgpu_right(
magma_int_t ngpu,
magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex_ptr d_lA[], magma_int_t ldda,
magma_int_t *info )
{
#define dlA(id, i, j) (d_lA[(id)] + (j) * ldda + (i))
#define dlP(id, i, j) (d_lP[(id)] + (j) * ldda + (i))
#define panel(j) (panel + (j))
#define tmppanel(j) (tmppanel + (j))
#define tmpprevpanel(j) (tmpprevpanel + (j))
#define STREAM_ID(i) (nqueue > 1 ? 1+((i)/nb)%(nqueue-1) : 0)
magmaFloatComplex z_one = MAGMA_C_MAKE( 1.0, 0.0 );
magmaFloatComplex mz_one = MAGMA_C_MAKE( -1.0, 0.0 );
float one = 1.0;
float m_one = -1.0;
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t j, nb, d, id, j_local, blkid, crosspoint, prevj, prevtrsmrows=0, nqueue = 5;
magmaFloatComplex *panel, *tmppanel0, *tmppanel1, *tmppanel, *tmpprevpanel;
magmaFloatComplex *d_lP[MagmaMaxGPUs], *dlpanel, *dlpanels[MagmaMaxGPUs];
magma_int_t rows, trsmrows, igpu, n_local[MagmaMaxGPUs], ldpanel;
magma_queue_t queues[MagmaMaxGPUs][10];
*info = 0;
if ( uplo != MagmaUpper && 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;
}
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
nb = magma_get_cpotrf_nb(n);
ldpanel = ldda;
magma_setdevice(0);
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &panel, 2 * nb * ldpanel )) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
tmppanel0 = panel;
tmppanel1 = tmppanel0 + nb * ldpanel;
if ((nb <= 1) || (nb >= n)) {
// Use unblocked code.
magma_cgetmatrix( n, n, dlA(0, 0, 0), ldda, panel, ldpanel);
lapackf77_cpotrf( uplo_, &n, panel, &ldpanel, info);
magma_csetmatrix( n, n, panel, ldpanel, dlA(0, 0, 0), ldda );
} else {
for( d = 0; d < ngpu; d++ ) {
// local-n and local-ld
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;
magma_setdevice(d);
magma_device_sync();
if (MAGMA_SUCCESS != magma_cmalloc( &d_lP[d], nb * ldda )) {
for( j = 0; j < d; j++ ) {
magma_setdevice(j);
magma_free( d_lP[d] );
}
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
for( j=0; j < nqueue; j++ ) {
magma_queue_create( &queues[d][j] );
}
}
//#define ENABLE_TIMER
#if defined (ENABLE_TIMER)
real_Double_t therk[4], tmtc, tcchol, tctrsm, tctm, tmnp, tcnp;
real_Double_t ttot_herk[4] = {0,0,0,0}, ttot_mtc = 0, ttot_cchol = 0, ttot_ctrsm = 0, ttot_ctm = 0, ttot_mnp = 0, ttot_cnp = 0;
printf("\n\n %10s %10s %10s %10s %10s %10s %10s %10s %10s %10s %10s %10s %10s %10s %10s\n",
"j", "nb", "row", "mtc", "CPU_np", "panel", "ctrsm", "CH+TRSM", "CPU", "dsyrk[0]", "dsyrk[1]", "dsyrk[2]", "dsyrk[3]", "ctm P", "gpu_np");
printf(" ====================================================================================================\n");
#endif
// Use blocked code.
if (uplo == MagmaUpper) {
printf( " === not supported, yet ===\n" );
} else {
blkid = -1;
if (ngpu == 4)
crosspoint = n;
else if (ngpu == 3)
crosspoint = n;
else if (ngpu == 2)
crosspoint = 20160;
else
crosspoint = 0;
crosspoint = 0; //n; //n -- > gpu always does next panel, 0 --> cpu always does next panel
crosspoint = n;
#if defined (ENABLE_TIMER)
real_Double_t tget = magma_wtime(), tset = 0.0, ttot = 0.0;
#endif
if ( n > nb ) {
// send first panel to cpu
magma_setdevice(0);
tmppanel = tmppanel0;
magma_cgetmatrix_async(n, nb,
dlA(0, 0, 0), ldda,
tmppanel(0), ldpanel,
queues[0][0] );
}
#if defined (ENABLE_TIMER)
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_device_sync();
}
tget = magma_wtime()-tget;
#endif
// Compute the Cholesky factorization A = L*L'
for (j = 0; (j + nb) < n; j += nb) {
#if defined (ENABLE_TIMER)
therk[0] = therk[1] = therk[2] = therk[3] = tmtc = tcchol = tctrsm = tctm = tmnp = tcnp = 0.0;
#endif
blkid += 1;
tmppanel = (blkid % 2 == 0) ? tmppanel0 : tmppanel1;
// Set the gpu number that holds the current panel
id = (j / nb) % ngpu;
magma_setdevice(id);
// Set the local index where the current panel is
j_local = j / (nb * ngpu) * nb;
rows = n - j;
// Wait for the panel on cpu
magma_queue_sync( queues[id][0] );
if (j > 0 && prevtrsmrows > crosspoint) {
#if defined (ENABLE_TIMER)
tcnp = magma_wtime();
#endif
tmpprevpanel = ((blkid - 1) % 2) == 0 ? tmppanel0 : tmppanel1;
blasf77_cgemm( MagmaNoTransStr, MagmaConjTransStr,
&rows, &nb, &nb,
&mz_one, tmpprevpanel(j), &ldpanel,
tmpprevpanel(j), &ldpanel,
&z_one, tmppanel(j), &ldpanel );
#if defined (ENABLE_TIMER)
tcnp = magma_wtime() - tcnp;
ttot_cnp += tcnp;
#endif
}
#if defined (ENABLE_TIMER)
tcchol = magma_wtime();
#endif
lapackf77_cpotrf(MagmaLowerStr, &nb, tmppanel(j), &ldpanel, info);
if (*info != 0) {
*info = *info + j;
break;
}
#if defined (ENABLE_TIMER)
tcchol = magma_wtime() - tcchol;
ttot_cchol += tcchol;
tctrsm = magma_wtime();
#endif
trsmrows = rows - nb;
if (trsmrows > 0) {
blasf77_ctrsm(MagmaRightStr, MagmaLowerStr, MagmaConjTransStr, MagmaNonUnitStr,
&trsmrows, &nb,
&z_one, tmppanel(j), &ldpanel,
tmppanel(j + nb), &ldpanel);
}
#if defined (ENABLE_TIMER)
tctrsm = magma_wtime() - tctrsm;
ttot_ctrsm += tctrsm;
tctm = magma_wtime();
#endif
d = (id + 1) % ngpu;
// send current panel to gpus
for (igpu = 0; igpu < ngpu; igpu++, d = (d + 1) % ngpu ) {
magma_int_t myrows = 0;
magma_int_t row_offset = 0;
if ( d == id ) {
dlpanel = dlA(d, j, j_local);
myrows = rows;
row_offset = 0;
} else {
dlpanel = dlP(d, 0, 0);
myrows = trsmrows;
row_offset = nb;
}
if (myrows > 0) {
magma_setdevice(d);
magma_csetmatrix_async(myrows, nb,
tmppanel(j + row_offset), ldpanel,
dlpanel, ldda, queues[d][0] );
}
}
/* make sure panel is on GPUs */
d = (id + 1) % ngpu;
for (igpu = 0; igpu < ngpu; igpu++, d = (d + 1) % ngpu ) {
magma_setdevice(d);
magma_queue_sync( queues[d][0] );
}
#if defined (ENABLE_TIMER)
tctm = magma_wtime() - tctm;
ttot_ctm += tctm;
#endif
if ( (j + nb) < n) {
magma_int_t offset = 0;
magma_int_t row_offset = 0;
if (j + nb + nb < n) {
d = (id + 1) % ngpu;
magma_setdevice(d);
magma_int_t j_local2 = (j + nb) / (nb * ngpu) * nb;
if (trsmrows <= crosspoint) {
#if defined (ENABLE_TIMER)
tmnp = magma_wtime();
#endif
// do gemm on look ahead panel
if ( d == id ) {
dlpanel = dlA(d, j + nb, j_local);
} else {
dlpanel = dlP(d, 0, 0);
}
magmablasSetKernelStream( queues[d][STREAM_ID(j_local2)] );
#define CHERK_ON_DIAG
#ifdef CHERK_ON_DIAG
magma_cherk( MagmaLower, MagmaNoTrans,
nb, nb,
m_one, dlpanel, ldda,
one, dlA(d, j + nb, j_local2), ldda);
magma_cgemm( MagmaNoTrans, MagmaConjTrans,
trsmrows-nb, nb, nb,
mz_one, dlpanel+nb, ldda,
dlpanel, ldda,
z_one, dlA(d, j + nb +nb, j_local2), ldda);
#else
magma_cgemm( MagmaNoTrans, MagmaConjTrans,
trsmrows, nb, nb,
mz_one, dlpanel, ldda,
dlpanel, ldda,
z_one, dlA(d, j + nb, j_local2), ldda);
#endif
#if defined (ENABLE_TIMER)
magma_device_sync();
tmnp = magma_wtime() - tmnp;
ttot_mnp += tmnp;
#endif
}
// send next panel to cpu
magma_queue_sync( queues[d][STREAM_ID(j_local2)] ); // make sure lookahead is done
tmppanel = ((blkid+1) % 2 == 0) ? tmppanel0 : tmppanel1;
magma_cgetmatrix_async(rows-nb, nb,
dlA(d, j+nb, j_local2), ldda,
tmppanel(j+nb), ldpanel,
queues[d][0] );
tmppanel = (blkid % 2 == 0) ? tmppanel0 : tmppanel1;
offset = j + nb + nb;
row_offset = nb;
} else {
offset = j + nb;
row_offset = 0;
}
if (n - offset > 0) {
// syrk on multiple gpu
for (d = 0; d < ngpu; d++ ) {
if ( d == id ) {
dlpanels[d] = dlA(d, j + nb + row_offset, j_local);
} else {
dlpanels[d] = dlP(d, row_offset, 0);
}
}
#if defined (ENABLE_TIMER)
for( d=0; d < ngpu; d++ ) therk[d] = magma_wtime();
#endif
//magmablasSetKernelStream( queues[d] );
//magma_cherk(MagmaLower, MagmaNoTrans, n - offset, nb,
// m_one, dlpanel, ldda,
// one, &d_lA[d][offset + offset*ldda], ldda );
#ifdef CHERK_ON_DIAG
magma_cherk_mgpu
#else
magma_cherk_mgpu2
#endif
(ngpu, MagmaLower, MagmaNoTrans,
nb, n - offset, nb,
m_one, dlpanels, ldda, 0,
one, d_lA, ldda, offset,
nqueue, queues );
#if defined (ENABLE_TIMER)
for( d=0; d < ngpu; d++ ) {
magma_setdevice(d);
magma_device_sync();
therk[d] = magma_wtime() - therk[d];
ttot_herk[d] += therk[d];
}
#endif
}
prevtrsmrows = trsmrows;
prevj = j;
#if defined (ENABLE_TIMER)
ttot += (tcnp+tcchol+tctrsm+therk[0]+therk[1]+therk[2]+tctm+tmnp);
printf("%10d %10d %10d %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf(%d) %10.3lf\n",
j, nb, rows, tmtc,
tcnp, // gemm
tcchol, // potrf
tctrsm, // trsm
(tcchol + tctrsm),
(tmtc+tcnp+tcchol+tctrsm),
therk[0], therk[1], therk[2], therk[3], // syrk
tctm, // copy panel to GPU
tmnp, // lookahead on GPU
(id + 1) % ngpu,
(tcnp+tcchol+tctrsm+therk[0]+therk[1]+therk[2]+tctm+tmnp));
fflush(0);
#endif
}
}
for( d = 0; d < ngpu; d++ ) {
magma_setdevice(d);
for( id=0; id < nqueue; id++ ) {
magma_queue_sync( queues[d][id] );
}
}
#if defined (ENABLE_TIMER)
printf("\n%10d %10d %10d %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf(-) %10.3lf\n",
n, n, 0, ttot_mtc,
ttot_cnp, // gemm
ttot_cchol, // potrf
ttot_ctrsm, // trsm
(ttot_cchol + ttot_ctrsm),
(ttot_mtc+ttot_cnp+ttot_cchol+ttot_ctrsm),
ttot_herk[0], ttot_herk[1], ttot_herk[2], ttot_herk[3], // syrk
ttot_ctm, // copy panel to GPU
ttot_mnp, // lookahead on GPU
(ttot_cnp+ttot_cchol+ttot_ctrsm+ttot_herk[0]+ttot_herk[1]+ttot_herk[2]+ttot_ctm+ttot_mnp));
printf("%10d %10d %10d %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf %10.3lf(-) %10.3lf (ratio)\n",
n, n, 0, ttot_mtc/ttot,
ttot_cnp/ttot, // gemm
ttot_cchol/ttot, // potrf
ttot_ctrsm/ttot, // trsm
(ttot_cchol + ttot_ctrsm)/ttot,
(ttot_mtc+ttot_cnp+ttot_cchol+ttot_ctrsm)/ttot,
ttot_herk[0]/ttot, ttot_herk[1]/ttot, ttot_herk[2]/ttot, ttot_herk[3]/ttot, // syrk
ttot_ctm/ttot, // copy panel to GPU
ttot_mnp/ttot, // lookahead on GPU
(ttot_cnp+ttot_cchol+ttot_ctrsm+ttot_herk[0]+ttot_herk[1]+ttot_herk[2]+ttot_ctm+ttot_mnp)/ttot);
#endif
// cholesky for the last block
if (j < n && *info == 0) {
rows = n - j;
id = (j / nb) % ngpu;
// Set the local index where the current panel is
j_local = j / (nb * ngpu) * nb;
magma_setdevice(id);
#if defined (ENABLE_TIMER)
tset = magma_wtime();
#endif
magma_cgetmatrix(rows, rows, dlA(id, j, j_local), ldda, panel(j), ldpanel);
lapackf77_cpotrf(MagmaLowerStr, &rows, panel(j), &ldpanel, info);
magma_csetmatrix(rows, rows, panel(j), ldpanel, dlA(id, j, j_local), ldda);
#if defined (ENABLE_TIMER)
tset = magma_wtime() - tset;
#endif
}
#if defined (ENABLE_TIMER)
printf( " matrix_get,set: %10.3lf %10.3lf -> %10.3lf\n",tget,tset,ttot+tget+tset );
#endif
} // end of else not upper
// clean up
for( d = 0; d < ngpu; d++ ) {
magma_setdevice(d);
for( j=0; j < nqueue; j++ ) {
magma_queue_destroy( queues[d][j] );
}
magma_free( d_lP[d] );
}
} // end of not lapack
// free workspace
magma_free_pinned( panel );
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
return *info;
} /* magma_cpotrf_mgpu_right */
virtual void run()
{
blasf77_cgemm( lapack_trans_const(transA), lapack_trans_const(transB),
&m, &n, &k, &alpha, A, &lda, B, &ldb, &beta, C, &ldc );
}
extern "C" magma_int_t
magma_cgetrf_nopiv(magma_int_t *m, magma_int_t *n, cuFloatComplex *a,
magma_int_t *lda, magma_int_t *info)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
CGETRF_NOPIV computes an LU factorization of a general M-by-N
matrix A without pivoting.
The factorization has the form
A = L * U
where 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
=========
M (input) INTEGER
The number of rows of the matrix A. M >= 0.
N (input) INTEGER
The number of columns of the matrix A. N >= 0.
A (input/output) COMPLEX*16 array, 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.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,M).
INFO (output) 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 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.
===================================================================== */
cuFloatComplex c_one = MAGMA_C_ONE;
magma_int_t a_dim1, a_offset, min_mn, i__3, i__4;
cuFloatComplex z__1;
magma_int_t j, jb, nb, iinfo;
a_dim1 = *lda;
a_offset = 1 + a_dim1;
a -= a_offset;
/* Function Body */
*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;
}
/* Determine the block size for this environment. */
nb = 128;
min_mn = min(*m,*n);
if (nb <= 1 || nb >= min_mn)
{
/* Use unblocked code. */
magma_cgetf2_nopiv(m, n, &a[a_offset], lda, info);
}
else
{
/* Use blocked code. */
for (j = 1; j <= min_mn; j += nb)
{
/* Computing MIN */
i__3 = min_mn - j + 1;
jb = min(i__3,nb);
/* Factor diagonal and subdiagonal blocks and test for exact
singularity. */
i__3 = *m - j + 1;
//magma_cgetf2_nopiv(&i__3, &jb, &a[j + j * a_dim1], lda, &iinfo);
i__3 -= jb;
magma_cgetf2_nopiv(&jb, &jb, &a[j + j * a_dim1], lda, &iinfo);
blasf77_ctrsm("R", "U", "N", "N", &i__3, &jb, &c_one,
&a[j + j * a_dim1], lda,
&a[j + jb + j * a_dim1], lda);
/* Adjust INFO */
if (*info == 0 && iinfo > 0)
*info = iinfo + j - 1;
if (j + jb <= *n)
{
/* Compute block row of U. */
i__3 = *n - j - jb + 1;
blasf77_ctrsm("Left", "Lower", "No transpose", "Unit", &jb, &i__3,
&c_one, &a[j + j * a_dim1], lda, &a[j + (j+jb)*a_dim1], lda);
if (j + jb <= *m)
{
/* Update trailing submatrix. */
i__3 = *m - j - jb + 1;
i__4 = *n - j - jb + 1;
z__1 = MAGMA_C_NEG_ONE;
blasf77_cgemm("No transpose", "No transpose", &i__3, &i__4, &jb,
&z__1, &a[j + jb + j * a_dim1], lda,
&a[j + (j + jb) * a_dim1], lda, &c_one,
&a[j + jb + (j + jb) * a_dim1], lda);
}
}
}
}
return *info;
} /* magma_cgetrf_nopiv */
/* ////////////////////////////////////////////////////////////////////////////
-- 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;
}
int main( int argc, char** argv)
{
real_Double_t gflops, gpu_perf, cpu_perf, gpu_time, cpu_time;
float error, work[1];
int transA = MagmaNoTrans;
int transB = MagmaNoTrans;
float Cnorm;
magma_int_t istart = 1024;
magma_int_t iend = 8194;
magma_int_t M, M0 = 0;
magma_int_t N, N0 = 0;
magma_int_t K, K0 = 0;
magma_int_t i;
magma_int_t Am, An, Bm, Bn;
magma_int_t szeA, szeB, szeC;
magma_int_t lda, ldb, ldc, ldda, lddb, lddc;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magmaFloatComplex *h_A, *h_B, *h_C, *h_C2;
magmaFloatComplex_ptr d_A, d_B, d_C;
magmaFloatComplex mzone = MAGMA_C_NEG_ONE;
magmaFloatComplex alpha = MAGMA_C_MAKE( 0.29, -0.86 );
magmaFloatComplex beta = MAGMA_C_MAKE( -0.48, 0.38 );
if (argc != 1){
for(i=1; i<argc; i++){
if ( strcmp("-N", argv[i]) == 0 ){
N0 = atoi(argv[++i]);
}
else if ( strcmp("-M", argv[i]) == 0 ){
M0 = atoi(argv[++i]);
}
else if ( strcmp("-K", argv[i]) == 0 ){
K0 = atoi(argv[++i]);
}
else if (strcmp("-NN", argv[i])==0){
transA = transB = MagmaNoTrans;
}
else if (strcmp("-TT", argv[i])==0){
transA = transB = MagmaTrans;
}
else if (strcmp("-NT", argv[i])==0){
transA = MagmaNoTrans;
transB = MagmaTrans;
}
else if (strcmp("-TN", argv[i])==0){
transA = MagmaTrans;
transB = MagmaNoTrans;
}
#if defined(PRECISION_z) || defined(PRECISION_c)
else if (strcmp("-NC", argv[i])==0){
transA = MagmaNoTrans;
transB = MagmaConjTrans;
}
else if (strcmp("-TC", argv[i])==0){
transA = MagmaTrans;
transB = MagmaConjTrans;
}
else if (strcmp("-CN", argv[i])==0){
transA = MagmaConjTrans;
transB = MagmaNoTrans;
}
else if (strcmp("-CT", argv[i])==0){
transA = MagmaConjTrans;
transB = MagmaTrans;
}
else if (strcmp("-CC", argv[i])==0){
transA = transB = MagmaConjTrans;
}
#endif
}
}
if ( (M0 != 0) && (N0 != 0) && (K0 != 0) )
iend = istart + 1;
M = N = K = iend;
if ( M0 != 0 ) M = M0;
if ( N0 != 0 ) N = N0;
if ( K0 != 0 ) K = K0;
if( transA == MagmaNoTrans ) {
Am = M;
An = K;
} else {
Am = K;
An = M;
}
if( transB == MagmaNoTrans ) {
Bm = K;
Bn = N;
} else {
Bm = N;
Bn = K;
}
/* Initialize */
magma_queue_t queue;
magma_device_t device[ MagmaMaxGPUs ];
int num = 0;
magma_err_t err;
magma_init();
err = magma_get_devices( device, MagmaMaxGPUs, &num );
if ( err != 0 || num < 1 ) {
fprintf( stderr, "magma_get_devices failed: %d\n", err );
exit(-1);
}
err = magma_queue_create( device[0], &queue );
if ( err != 0 ) {
fprintf( stderr, "magma_queue_create failed: %d\n", err );
exit(-1);
}
lda = ldc = M;
ldb = Bm;
ldda = lddc = ((M+31)/32)*32;
lddb = ((ldb+31)/32)*32;
K+=32;
M+=32;
N +=32;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, lda*K );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, ldb*Bn );
TESTING_MALLOC_CPU( h_C, magmaFloatComplex, ldc*N );
TESTING_MALLOC_CPU( h_C2, magmaFloatComplex, ldc*N );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*K );
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, lddb*Bn );
TESTING_MALLOC_DEV( d_C, magmaFloatComplex, lddc*N );
printf("\nUsage: \n");
printf(" testing_cgemm [-NN|NT|TN|TT] [-N %d] \n\n", 1024);
printf("\n");
printf("Testing transA = %c transB = %c\n", transA, transB);
printf(" M N K clAmdBlas GFLop/s (sec) CPU GFlop/s (sec) error\n");
printf("===========================================================================\n");
for(i=istart; i<iend; i = (int)(i*1.25) )
{
M = N = K = i;
if ( M0 != 0 ) M = M0;
if ( N0 != 0 ) N = N0;
if ( K0 != 0 ) K = K0;
if( transA == MagmaNoTrans ) {
lda = Am = M;
An = K;
} else {
lda = Am = K;
An = M;
}
if( transB == MagmaNoTrans ) {
ldb = Bm = K;
Bn = N;
} else {
ldb = Bm = N;
Bn = K;
}
gflops = FLOPS( (float)M, (float)N, (float)K ) * 1e-9;
ldc = M;
ldda = ((lda+31)/32)*32;
lddb = ((ldb+31)/32)*32;
lddc = ((ldc+31)/32)*32;
szeA = lda * An;
szeB = ldb * Bn;
szeC = ldc * N;
/* Initialize the matrices */
lapackf77_clarnv( &ione, ISEED, &szeA, h_A );
lapackf77_clarnv( &ione, ISEED, &szeB, h_B );
lapackf77_clarnv( &ione, ISEED, &szeC, h_C );
/* =====================================================================
Performs operation using MAGMA-BLAS
=================================================================== */
magma_csetmatrix( Am, An, h_A, 0, lda, d_A, 0, ldda, queue );
magma_csetmatrix( Bm, Bn, h_B, 0, ldb, d_B, 0, lddb, queue );
magma_csetmatrix( M, N, h_C, 0, ldc, d_C, 0, lddc, queue );
magma_cgemm( transA, transB, M, N, K,
alpha, d_A, 0, ldda,
d_B, 0, lddb,
beta, d_C, 0, lddc, queue );
magma_csetmatrix( M, N, h_C, 0, ldc, d_C, 0, lddc, queue );
magma_queue_sync( queue );
gpu_time = magma_wtime();
magma_cgemm( transA, transB, M, N, K,
alpha, d_A, 0, ldda,
d_B, 0, lddb,
beta, d_C, 0, lddc, queue );
magma_queue_sync( queue);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
magma_cgetmatrix( M, N, d_C, 0, lddc, h_C2, 0, ldc, queue );
/* =====================================================================
Performs operation using CPU-BLAS
=================================================================== */
cpu_time = magma_wtime();
blasf77_cgemm( lapack_const(transA), lapack_const(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;
// |C_magma - C_lapack| / |C_lapack|
Cnorm = lapackf77_clange( "M", &M, &N, h_C, &ldc, work );
/* =====================================================================
Error Computation and Performance Compariosn
=================================================================== */
blasf77_caxpy(&szeC, &mzone, h_C, &ione, h_C2, &ione);
error = lapackf77_clange("M", &M, &N, h_C2, &ldc, work)/Cnorm;
printf("%5d %5d %5d %8.2f (%6.2f) %6.2f (%6.2f) %e\n",
M, N, K, gpu_perf, gpu_time, cpu_perf, cpu_time, error);
}
/* Memory clean up */
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_C );
TESTING_FREE_CPU( h_C2 );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
TESTING_FREE_DEV( d_C );
magma_queue_destroy( queue );
magma_finalize();
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cgels
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
float gpu_error, cpu_error, error, Anorm, work[1];
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_A2, *h_B, *h_X, *h_R, *tau, *h_work, tmp[1];
magmaFloatComplex *d_A, *d_B;
magma_int_t M, N, size, nrhs, lda, ldb, ldda, lddb, min_mn, max_mn, nb, info;
magma_int_t lworkgpu, lhwork, lhwork2;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_int_t status = 0;
float tol = opts.tolerance * lapackf77_slamch("E");
nrhs = opts.nrhs;
printf(" ||b-Ax|| / (N||A||) ||dx-x||/(N||A||)\n");
printf(" M N NRHS CPU GFlop/s (sec) GPU GFlop/s (sec) CPU GPU \n");
printf("===================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
if ( M < N ) {
printf( "%5d %5d %5d skipping because M < N is not yet supported.\n", (int) M, (int) N, (int) nrhs );
continue;
}
min_mn = min(M, N);
max_mn = max(M, N);
lda = M;
ldb = max_mn;
size = lda*N;
ldda = ((M+31)/32)*32;
lddb = ((max_mn+31)/32)*32;
nb = magma_get_cgeqrf_nb(M);
gflops = (FLOPS_CGEQRF( M, N ) + FLOPS_CGEQRS( M, N, nrhs )) / 1e9;
lworkgpu = (M - N + nb)*(nrhs + nb) + nrhs*nb;
// query for workspace size
lhwork = -1;
lapackf77_cgeqrf(&M, &N, NULL, &M, NULL, tmp, &lhwork, &info);
lhwork2 = (magma_int_t) MAGMA_C_REAL( tmp[0] );
lhwork = -1;
lapackf77_cunmqr( MagmaLeftStr, MagmaConjTransStr,
&M, &nrhs, &min_mn, NULL, &lda, NULL,
NULL, &ldb, tmp, &lhwork, &info);
lhwork = (magma_int_t) MAGMA_C_REAL( tmp[0] );
lhwork = max( max( lhwork, lhwork2 ), lworkgpu );
TESTING_MALLOC_CPU( tau, magmaFloatComplex, min_mn );
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( h_A2, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, ldb*nrhs );
TESTING_MALLOC_CPU( h_X, magmaFloatComplex, ldb*nrhs );
TESTING_MALLOC_CPU( h_R, magmaFloatComplex, ldb*nrhs );
TESTING_MALLOC_CPU( h_work, magmaFloatComplex, lhwork );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, lddb*nrhs );
/* Initialize the matrices */
lapackf77_clarnv( &ione, ISEED, &size, h_A );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_A2, &lda );
// make random RHS
size = M*nrhs;
lapackf77_clarnv( &ione, ISEED, &size, h_B );
lapackf77_clacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_R, &ldb );
// make consistent RHS
//size = N*nrhs;
//lapackf77_clarnv( &ione, ISEED, &size, h_X );
//blasf77_cgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
// &c_one, h_A, &lda,
// h_X, &ldb,
// &c_zero, h_B, &ldb );
//lapackf77_clacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_R, &ldb );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_csetmatrix( M, N, h_A, lda, d_A, ldda );
magma_csetmatrix( M, nrhs, h_B, ldb, d_B, lddb );
gpu_time = magma_wtime();
magma_cgels3_gpu( MagmaNoTrans, M, N, nrhs, d_A, ldda,
d_B, lddb, h_work, lworkgpu, &info);
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cgels3_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// Get the solution in h_X
magma_cgetmatrix( N, nrhs, d_B, lddb, h_X, ldb );
// compute the residual
blasf77_cgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
&c_neg_one, h_A, &lda,
h_X, &ldb,
&c_one, h_R, &ldb);
Anorm = lapackf77_clange("f", &M, &N, h_A, &lda, work);
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
lapackf77_clacpy( MagmaUpperLowerStr, &M, &nrhs, h_B, &ldb, h_X, &ldb );
cpu_time = magma_wtime();
lapackf77_cgels( MagmaNoTransStr, &M, &N, &nrhs,
h_A, &lda, h_X, &ldb, h_work, &lhwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cgels returned error %d: %s.\n",
(int) info, magma_strerror( info ));
blasf77_cgemm( MagmaNoTransStr, MagmaNoTransStr, &M, &nrhs, &N,
&c_neg_one, h_A2, &lda,
h_X, &ldb,
&c_one, h_B, &ldb);
cpu_error = lapackf77_clange("f", &M, &nrhs, h_B, &ldb, work) / (min_mn*Anorm);
gpu_error = lapackf77_clange("f", &M, &nrhs, h_R, &ldb, work) / (min_mn*Anorm);
// error relative to LAPACK
size = M*nrhs;
blasf77_caxpy( &size, &c_neg_one, h_B, &ione, h_R, &ione );
error = lapackf77_clange("f", &M, &nrhs, h_R, &ldb, work) / (min_mn*Anorm);
printf("%5d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %8.2e",
(int) M, (int) N, (int) nrhs,
cpu_perf, cpu_time, gpu_perf, gpu_time, cpu_error, gpu_error, error );
if ( M == N ) {
printf( " %s\n", (gpu_error < tol && error < tol ? "ok" : "failed"));
status += ! (gpu_error < tol && error < tol);
}
else {
printf( " %s\n", (error < tol ? "ok" : "failed"));
status += ! (error < tol);
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_A2 );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_X );
TESTING_FREE_CPU( h_R );
TESTING_FREE_CPU( h_work );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
int main( int argc, char** argv )
{
magma_init();
cublasHandle_t handle;
cudaSetDevice( 0 );
cublasCreate( &handle );
magmaFloatComplex *A, *B, *C;
magmaFloatComplex *dA, *dB, *dC;
float error, work[1];
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magma_int_t ione = 1;
magma_int_t ISEED[4] = { 1, 2, 3, 4 };
magma_int_t n, lda, ldda, size, info;
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf(" N |dC - C|/|C|\n");
printf("====================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
// for this simple case, all matrices are N-by-N
n = opts.nsize[itest];
lda = n;
ldda = ((n+31)/32)*32;
magma_cmalloc_cpu( &A, lda*n );
magma_cmalloc_cpu( &B, lda*n );
magma_cmalloc_cpu( &C, lda*n );
magma_cmalloc( &dA, ldda*n );
magma_cmalloc( &dB, ldda*n );
magma_cmalloc( &dC, ldda*n );
// initialize matrices
size = lda*n;
lapackf77_clarnv( &ione, ISEED, &size, A );
lapackf77_clarnv( &ione, ISEED, &size, B );
lapackf77_clarnv( &ione, ISEED, &size, C );
// increase diagonal to be SPD
for( int i=0; i < n; ++i ) {
C[i+i*lda] = MAGMA_C_ADD( C[i+i*lda], MAGMA_C_MAKE( n*n, 0 ));
}
magma_csetmatrix( n, n, A, lda, dA, ldda );
magma_csetmatrix( n, n, B, lda, dB, ldda );
magma_csetmatrix( n, n, C, lda, dC, ldda );
// compute with cublas
cublasCgemm( handle, CUBLAS_OP_N, CUBLAS_OP_N, n, n, n,
&c_neg_one, dA, ldda, dB, ldda, &c_one, dC, ldda );
magma_cpotrf_gpu( MagmaLower, n, dC, ldda, &info );
if (info != 0)
printf("magma_cpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// compute with LAPACK
blasf77_cgemm( MagmaNoTransStr, MagmaNoTransStr, &n, &n, &n,
&c_neg_one, A, &lda, B, &lda, &c_one, C, &lda );
lapackf77_cpotrf( MagmaLowerStr, &n, C, &lda, &info );
if (info != 0)
printf("lapackf77_cpotrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
// compute difference, |dC - C| / |C|
magma_cgetmatrix( n, n, dC, ldda, A, lda );
blasf77_caxpy( &size, &c_neg_one, C, &ione, A, &ione );
error = lapackf77_clange( "F", &n, &n, A, &lda, work )
/ lapackf77_clange( "F", &n, &n, C, &lda, work );
printf( "%5d %8.2e %s\n",
(int) n, error, (error < tol ? "ok" : "failed"));
status += ! (error < tol);
magma_free( dA );
magma_free( dB );
magma_free( dC );
magma_free_cpu( A );
magma_free_cpu( B );
magma_free_cpu( C );
fflush( stdout );
}
}
cublasDestroy( handle );
magma_finalize();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing cposv_gpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, cpu_perf, cpu_time, gpu_perf, gpu_time;
float error, Rnorm, Anorm, Xnorm, *work;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex *h_A, *h_B, *h_X;
magmaFloatComplex_ptr d_A, d_B;
magma_int_t N, lda, ldb, ldda, lddb, info, sizeA, sizeB;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
float tol = opts.tolerance * lapackf77_slamch("E");
printf("uplo = %s\n", lapack_uplo_const(opts.uplo) );
printf(" N NRHS CPU Gflop/s (sec) GPU GFlop/s (sec) ||B - AX|| / N*||A||*||X||\n");
printf("================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = ldb = N;
ldda = ((N+31)/32)*32;
lddb = ldda;
gflops = ( FLOPS_CPOTRF( N ) + FLOPS_CPOTRS( N, opts.nrhs ) ) / 1e9;
TESTING_MALLOC_CPU( h_A, magmaFloatComplex, lda*N );
TESTING_MALLOC_CPU( h_B, magmaFloatComplex, ldb*opts.nrhs );
TESTING_MALLOC_CPU( h_X, magmaFloatComplex, ldb*opts.nrhs );
TESTING_MALLOC_CPU( work, float, N );
TESTING_MALLOC_DEV( d_A, magmaFloatComplex, ldda*N );
TESTING_MALLOC_DEV( d_B, magmaFloatComplex, lddb*opts.nrhs );
/* ====================================================================
Initialize the matrix
=================================================================== */
sizeA = lda*N;
sizeB = ldb*opts.nrhs;
lapackf77_clarnv( &ione, ISEED, &sizeA, h_A );
lapackf77_clarnv( &ione, ISEED, &sizeB, h_B );
magma_cmake_hpd( N, h_A, lda );
magma_csetmatrix( N, N, h_A, N, d_A, ldda );
magma_csetmatrix( N, opts.nrhs, h_B, N, d_B, lddb );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
magma_cposv_gpu( opts.uplo, N, opts.nrhs, d_A, ldda, d_B, lddb, &info );
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_cpotrf_gpu returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Residual
=================================================================== */
magma_cgetmatrix( N, opts.nrhs, d_B, lddb, h_X, ldb );
Anorm = lapackf77_clange("I", &N, &N, h_A, &lda, work);
Xnorm = lapackf77_clange("I", &N, &opts.nrhs, h_X, &ldb, work);
blasf77_cgemm( MagmaNoTransStr, MagmaNoTransStr, &N, &opts.nrhs, &N,
&c_one, h_A, &lda,
h_X, &ldb,
&c_neg_one, h_B, &ldb );
Rnorm = lapackf77_clange("I", &N, &opts.nrhs, h_B, &ldb, work);
error = Rnorm/(N*Anorm*Xnorm);
status += ! (error < tol);
/* ====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_cposv( lapack_uplo_const(opts.uplo), &N, &opts.nrhs, h_A, &lda, h_B, &ldb, &info );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_cposv returned error %d: %s.\n",
(int) info, magma_strerror( info ));
printf( "%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %s\n",
(int) N, (int) opts.nrhs, cpu_perf, cpu_time, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed"));
}
else {
printf( "%5d %5d --- ( --- ) %7.2f (%7.2f) %8.2e %s\n",
(int) N, (int) opts.nrhs, gpu_perf, gpu_time,
error, (error < tol ? "ok" : "failed"));
}
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_B );
TESTING_FREE_CPU( h_X );
TESTING_FREE_CPU( work );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_B );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing chegvd
*/
int main( int argc, char** argv)
{
TESTING_CUDA_INIT();
cuFloatComplex *h_A, *h_R, *h_B, *h_S, *h_work;
float *rwork, *w1, *w2;
magma_int_t *iwork;
float gpu_time, cpu_time;
magma_timestr_t start, end;
/* Matrix size */
magma_int_t N=0, n2;
magma_int_t size[4] = {1024,2048,4100,6001};
magma_int_t i, itype, info;
magma_int_t ione = 1, izero = 0;
magma_int_t five = 5;
cuFloatComplex c_zero = MAGMA_C_ZERO;
cuFloatComplex c_one = MAGMA_C_ONE;
cuFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
float d_one = 1.;
float d_neg_one = -1.;
float d_ten = 10.;
magma_int_t ISEED[4] = {0,0,0,1};
const char *uplo = MagmaLowerStr;
const char *jobz = MagmaVectorsStr;
itype = 1;
magma_int_t checkres;
float result[4];
int flagN = 0;
if (argc != 1){
for(i = 1; i<argc; i++){
if (strcmp("-N", argv[i])==0){
N = atoi(argv[++i]);
if (N>0){
printf(" testing_chegvd -N %d\n\n", (int) N);
flagN=1;
}
else {
printf("\nUsage: \n");
printf(" testing_chegvd -N %d\n\n", (int) N);
exit(1);
}
}
if (strcmp("-itype", argv[i])==0){
itype = atoi(argv[++i]);
if (itype>0 && itype <= 3){
printf(" testing_chegvd -itype %d\n\n", (int) itype);
}
else {
printf("\nUsage: \n");
printf(" testing_chegvd -itype %d\n\n", (int) itype);
exit(1);
}
}
if (strcmp("-L", argv[i])==0){
uplo = MagmaLowerStr;
printf(" testing_chegvd -L");
}
if (strcmp("-U", argv[i])==0){
uplo = MagmaUpperStr;
printf(" testing_chegvd -U");
}
}
} else {
printf("\nUsage: \n");
printf(" testing_chegvd -L/U -N %d -itype %d\n\n", 1024, 1);
}
if(!flagN)
N = size[3];
checkres = getenv("MAGMA_TESTINGS_CHECK") != NULL;
n2 = N * N;
/* Allocate host memory for the matrix */
TESTING_MALLOC( h_A, cuFloatComplex, n2);
TESTING_MALLOC( h_B, cuFloatComplex, n2);
TESTING_MALLOC( w1, float , N);
TESTING_MALLOC( w2, float , N);
TESTING_HOSTALLOC(h_R, cuFloatComplex, n2);
TESTING_HOSTALLOC(h_S, cuFloatComplex, n2);
magma_int_t nb = magma_get_chetrd_nb(N);
magma_int_t lwork = 2*N*nb + N*N;
magma_int_t lrwork = 1 + 5*N +2*N*N;
magma_int_t liwork = 3 + 5*N;
TESTING_HOSTALLOC(h_work, cuFloatComplex, lwork);
TESTING_MALLOC( rwork, float, lrwork);
TESTING_MALLOC( iwork, magma_int_t, liwork);
printf(" N CPU Time(s) GPU Time(s) \n");
printf("===================================\n");
for(i=0; i<4; i++){
if (!flagN){
N = size[i];
n2 = N*N;
}
/* Initialize the matrix */
lapackf77_clarnv( &ione, ISEED, &n2, h_A );
//lapackf77_clatms( &N, &N, "U", ISEED, "P", w1, &five, &d_ten,
// &d_one, &N, &N, uplo, h_B, &N, h_work, &info);
//lapackf77_claset( "A", &N, &N, &c_zero, &c_one, h_B, &N);
lapackf77_clarnv( &ione, ISEED, &n2, h_B );
/* increase the diagonal */
{
magma_int_t i, j;
for(i=0; i<N; i++) {
MAGMA_C_SET2REAL( h_B[i*N+i], MAGMA_C_REAL(h_B[i*N+i]) + 1.*N );
MAGMA_C_SET2REAL( h_A[i*N+i], MAGMA_C_REAL(h_A[i*N+i]) );
}
}
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
magma_chegvd(itype, jobz[0], uplo[0],
N, h_R, N, h_S, N, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
start = get_current_time();
magma_chegvd(itype, jobz[0], uplo[0],
N, h_R, N, h_S, N, w1,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
end = get_current_time();
gpu_time = GetTimerValue(start,end)/1000.;
if ( checkres ) {
/* =====================================================================
Check the results following the LAPACK's [zc]hegvd routine.
A x = lambda B x is solved
and the following 3 tests computed:
(1) | A Z - B Z D | / ( |A||Z| N ) (itype = 1)
| A B Z - Z D | / ( |A||Z| N ) (itype = 2)
| B A Z - Z D | / ( |A||Z| N ) (itype = 3)
(2) | I - V V' B | / ( N ) (itype = 1,2)
| B - V V' | / ( |B| N ) (itype = 3)
(3) | S(with V) - S(w/o V) | / | S |
=================================================================== */
float temp1, temp2;
cuFloatComplex *tau;
if (itype == 1 || itype == 2){
lapackf77_claset( "A", &N, &N, &c_zero, &c_one, h_S, &N);
blasf77_cgemm("N", "C", &N, &N, &N, &c_one, h_R, &N, h_R, &N, &c_zero, h_work, &N);
blasf77_chemm("R", uplo, &N, &N, &c_neg_one, h_B, &N, h_work, &N, &c_one, h_S, &N);
result[1]= lapackf77_clange("1", &N, &N, h_S, &N, rwork) / N;
}
else if (itype == 3){
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N);
blasf77_cherk(uplo, "N", &N, &N, &d_neg_one, h_R, &N, &d_one, h_S, &N);
result[1]= lapackf77_clanhe("1",uplo, &N, h_S, &N, rwork) / N / lapackf77_clanhe("1",uplo, &N, h_B, &N, rwork);
}
result[0] = 1.;
result[0] /= lapackf77_clanhe("1",uplo, &N, h_A, &N, rwork);
result[0] /= lapackf77_clange("1",&N , &N, h_R, &N, rwork);
if (itype == 1){
blasf77_chemm("L", uplo, &N, &N, &c_one, h_A, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i<N; ++i)
blasf77_csscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_chemm("L", uplo, &N, &N, &c_neg_one, h_B, &N, h_R, &N, &c_one, h_work, &N);
result[0] *= lapackf77_clange("1", &N, &N, h_work, &N, rwork)/N;
}
else if (itype == 2){
blasf77_chemm("L", uplo, &N, &N, &c_one, h_B, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i<N; ++i)
blasf77_csscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_chemm("L", uplo, &N, &N, &c_one, h_A, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_clange("1", &N, &N, h_R, &N, rwork)/N;
}
else if (itype == 3){
blasf77_chemm("L", uplo, &N, &N, &c_one, h_A, &N, h_R, &N, &c_zero, h_work, &N);
for(int i=0; i<N; ++i)
blasf77_csscal(&N, &w1[i], &h_R[i*N], &ione);
blasf77_chemm("L", uplo, &N, &N, &c_one, h_B, &N, h_work, &N, &c_neg_one, h_R, &N);
result[0] *= lapackf77_clange("1", &N, &N, h_R, &N, rwork)/N;
}
/* lapackf77_chet21(&ione, uplo, &N, &izero,
h_A, &N,
w1, w1,
h_R, &N,
h_R, &N,
tau, h_work, rwork, &result[0]);
*/
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_A, &N, h_R, &N );
lapackf77_clacpy( MagmaUpperLowerStr, &N, &N, h_B, &N, h_S, &N );
magma_chegvd(itype, 'N', uplo[0],
N, h_R, N, h_S, N, w2,
h_work, lwork,
rwork, lrwork,
iwork, liwork,
&info);
temp1 = temp2 = 0;
for(int j=0; j<N; j++){
temp1 = max(temp1, absv(w1[j]));
temp1 = max(temp1, absv(w2[j]));
temp2 = max(temp2, absv(w1[j]-w2[j]));
}
result[2] = temp2 / temp1;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
start = get_current_time();
lapackf77_chegvd(&itype, jobz, uplo,
&N, h_A, &N, h_B, &N, w2,
h_work, &lwork,
rwork, &lrwork,
iwork, &liwork,
&info);
end = get_current_time();
if (info < 0)
printf("Argument %d of chegvd had an illegal value.\n", (int) -info);
cpu_time = GetTimerValue(start,end)/1000.;
/* =====================================================================
Print execution time
=================================================================== */
printf("%5d %6.2f %6.2f\n",
(int) N, cpu_time, gpu_time);
if ( checkres ){
printf("Testing the eigenvalues and eigenvectors for correctness:\n");
if(itype==1)
printf("(1) | A Z - B Z D | / (|A| |Z| N) = %e\n", result[0]);
else if(itype==2)
printf("(1) | A B Z - Z D | / (|A| |Z| N) = %e\n", result[0]);
else if(itype==3)
printf("(1) | B A Z - Z D | / (|A| |Z| N) = %e\n", result[0]);
if(itype==1 || itype ==2)
printf("(2) | I - Z Z' B | / N = %e\n", result[1]);
else
printf("(2) | B - Z Z' | / (|B| N) = %e\n", result[1]);
printf("(3) | D(w/ Z)-D(w/o Z)|/ |D| = %e\n\n", result[2]);
}
if (flagN)
break;
}
/* Memory clean up */
TESTING_FREE( h_A);
TESTING_FREE( h_B);
TESTING_FREE( w1);
TESTING_FREE( w2);
TESTING_FREE( rwork);
TESTING_FREE( iwork);
TESTING_HOSTFREE(h_work);
TESTING_HOSTFREE( h_R);
TESTING_HOSTFREE( h_S);
/* Shutdown */
TESTING_CUDA_FINALIZE();
}