K* malloc(size_t size) {
K* ptr;
if(MAGMA_SUCCESS !=
magma_malloc_cpu( (void**)&ptr, size*sizeof(K) )) {
std::cerr << "MAGMA ERROR: malloc failed\n";
exit(-1);
}
return ptr;
}
extern "C" magma_int_t
magma_zbulge_back_m(
magma_int_t ngpu,
magma_uplo_t uplo,
magma_int_t n, magma_int_t nb,
magma_int_t ne, magma_int_t Vblksiz,
magmaDoubleComplex *Z, magma_int_t ldz,
magmaDoubleComplex *V, magma_int_t ldv,
magmaDoubleComplex *TAU,
magmaDoubleComplex *T, magma_int_t ldt,
magma_int_t* info)
{
magma_int_t threads = magma_get_parallel_numthreads();
magma_int_t mklth = magma_get_lapack_numthreads();
magma_set_lapack_numthreads(1);
real_Double_t timeaplQ2=0.0;
double f= 1.;
magma_int_t n_gpu = ne;
//#if defined(PRECISION_s) || defined(PRECISION_d)
// double gpu_cpu_perf = 32; //gpu over cpu performance
//#else
// double gpu_cpu_perf = 32; // gpu over cpu performance
//#endif
double perf_temp= .85;
double perf_temp2= perf_temp;
for (magma_int_t itmp=1; itmp < ngpu; ++itmp)
perf_temp2 *= perf_temp;
magma_int_t gpu_cpu_perf = magma_get_zbulge_gcperf();
if (threads > 1) {
f = 1. / (1. + (double)(threads-1)/ ((double)gpu_cpu_perf*(1.-perf_temp2)/(1.-perf_temp)));
n_gpu = (magma_int_t)(f*ne);
}
/****************************************************
* apply V2 from left to the eigenvectors Z. dZ = (I-V2*T2*V2')*Z
* **************************************************/
timeaplQ2 = magma_wtime();
/*============================
* use GPU+CPU's
*==========================*/
//n_gpu = ne;
if (n_gpu < ne) {
// define the size of Q to be done on CPU's and the size on GPU's
// note that GPU use Q(1:N_GPU) and CPU use Q(N_GPU+1:N)
#ifdef ENABLE_DEBUG
printf("---> calling GPU + CPU(if N_CPU > 0) to apply V2 to Z with NE %d N_GPU %d N_CPU %d\n",ne, n_gpu, ne-n_gpu);
#endif
magma_zapplyQ_m_data data_applyQ(ngpu, threads, n, ne, n_gpu, nb, Vblksiz, Z, ldz, V, ldv, TAU, T, ldt);
magma_zapplyQ_m_id_data* arg;
magma_malloc_cpu((void**) &arg, threads*sizeof(magma_zapplyQ_m_id_data));
pthread_t* thread_id;
magma_malloc_cpu((void**) &thread_id, threads*sizeof(pthread_t));
pthread_attr_t thread_attr;
// ===============================
// relaunch thread to apply Q
// ===============================
// Set one thread per core
pthread_attr_init(&thread_attr);
pthread_attr_setscope(&thread_attr, PTHREAD_SCOPE_SYSTEM);
pthread_setconcurrency(threads);
// Launch threads
for (magma_int_t thread = 1; thread < threads; thread++) {
arg[thread] = magma_zapplyQ_m_id_data(thread, &data_applyQ);
pthread_create(&thread_id[thread], &thread_attr, magma_zapplyQ_m_parallel_section, &arg[thread]);
}
arg[0] = magma_zapplyQ_m_id_data(0, &data_applyQ);
magma_zapplyQ_m_parallel_section(&arg[0]);
// Wait for completion
for (magma_int_t thread = 1; thread < threads; thread++) {
void *exitcodep;
pthread_join(thread_id[thread], &exitcodep);
}
magma_free_cpu(thread_id);
magma_free_cpu(arg);
/*============================
* use only GPU
*==========================*/
} else {
magma_zbulge_applyQ_v2_m(ngpu, MagmaLeft, ne, n, nb, Vblksiz, Z, ldz, V, ldv, T, ldt, info);
magma_device_sync();
}
timeaplQ2 = magma_wtime()-timeaplQ2;
magma_set_lapack_numthreads(mklth);
return MAGMA_SUCCESS;
}
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, t1, t2;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_trans_t trans[] = { MagmaNoTrans, MagmaConjTrans, MagmaTrans };
magma_uplo_t uplo [] = { MagmaLower, MagmaUpper };
magma_diag_t diag [] = { MagmaUnit, MagmaNonUnit };
magma_side_t side [] = { MagmaLeft, MagmaRight };
magmaDoubleComplex *A, *B, *C, *C2, *LU;
magmaDoubleComplex *dA, *dB, *dC1, *dC2;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 0.5, 0.1 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( 0.7, 0.2 );
double dalpha = 0.6;
double dbeta = 0.8;
double work[1], error, total_error;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t m, n, k, size, maxn, ld, info;
magma_int_t *piv;
magma_int_t err;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf( "Compares magma wrapper function to cublas function; all diffs should be exactly 0.\n\n" );
total_error = 0.;
for( int itest = 0; itest < opts.ntest; ++itest ) {
m = opts.msize[itest];
n = opts.nsize[itest];
k = opts.ksize[itest];
printf("=========================================================================\n");
printf( "m=%d, n=%d, k=%d\n", (int) m, (int) n, (int) k );
// allocate matrices
// over-allocate so they can be any combination of {m,n,k} x {m,n,k}.
maxn = max( max( m, n ), k );
ld = max( 1, maxn );
size = ld*maxn;
err = magma_malloc_cpu( (void**) &piv, maxn*sizeof(magma_int_t) ); assert( err == 0 );
err = magma_zmalloc_pinned( &A, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &B, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &C, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &C2, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &LU, size ); assert( err == 0 );
err = magma_zmalloc( &dA, size ); assert( err == 0 );
err = magma_zmalloc( &dB, size ); assert( err == 0 );
err = magma_zmalloc( &dC1, size ); assert( err == 0 );
err = magma_zmalloc( &dC2, size ); assert( err == 0 );
// initialize matrices
size = maxn*maxn;
lapackf77_zlarnv( &ione, ISEED, &size, A );
lapackf77_zlarnv( &ione, ISEED, &size, B );
lapackf77_zlarnv( &ione, ISEED, &size, C );
printf( "========== Level 1 BLAS ==========\n" );
// ----- test ZSWAP
// swap columns 2 and 3 of dA, then copy to C2 and compare with A
if ( n >= 3 ) {
magma_zsetmatrix( m, n, A, ld, dA, ld );
magma_zsetmatrix( m, n, A, ld, dB, ld );
magma_zswap( m, dA(0,1), 1, dA(0,2), 1 );
magma_zswap( m, dB(0,1), 1, dB(0,2), 1 );
// check results, storing diff between magma and cuda calls in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dA, 1, dB, 1 );
magma_zgetmatrix( m, n, dB, ld, C2, ld );
error = lapackf77_zlange( "F", &m, &k, C2, &ld, work );
total_error += error;
printf( "zswap diff %.2g\n", error );
}
else {
printf( "zswap skipped for n < 3\n" );
}
// ----- test IZAMAX
// get argmax of column of A
magma_zsetmatrix( m, k, A, ld, dA, ld );
error = 0;
for( int j = 0; j < k; ++j ) {
magma_int_t i1 = magma_izamax( m, dA(0,j), 1 );
int i2; // NOT magma_int_t, for cublas
cublasIzamax( handle, m, dA(0,j), 1, &i2 );
// todo need sync here?
assert( i1 == i2 );
error += abs( i1 - i2 );
}
total_error += error;
gflops = (double)m * k / 1e9;
printf( "izamax diff %.2g\n", error );
printf( "\n" );
printf( "========== Level 2 BLAS ==========\n" );
// ----- test ZGEMV
// c = alpha*A*b + beta*c, with A m*n; b,c m or n-vectors
// try no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
magma_zsetmatrix( m, n, A, ld, dA, ld );
magma_zsetvector( maxn, B, 1, dB, 1 );
magma_zsetvector( maxn, C, 1, dC1, 1 );
magma_zsetvector( maxn, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_zgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZgemv( handle, cublas_trans_const(trans[ia]),
m, n, &alpha, dA, ld, dB, 1, &beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
size = (trans[ia] == MagmaNoTrans ? m : n);
cublasZaxpy( handle, size, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetvector( size, dC2, 1, C2, 1 );
error = lapackf77_zlange( "F", &size, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZGEMV( m, n ) / 1e9;
printf( "zgemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_trans_const(trans[ia]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test ZHEMV
// c = alpha*A*b + beta*c, with A m*m symmetric; b,c m-vectors
// try upper/lower
for( int iu = 0; iu < 2; ++iu ) {
magma_zsetmatrix( m, m, A, ld, dA, ld );
magma_zsetvector( m, B, 1, dB, 1 );
magma_zsetvector( m, C, 1, dC1, 1 );
magma_zsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_zhemv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZhemv( handle, cublas_uplo_const(uplo[iu]),
m, &alpha, dA, ld, dB, 1, &beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_zlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHEMV( m ) / 1e9;
printf( "zhemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test ZTRSV
// solve A*c = c, with A m*m triangular; c m-vector
// try upper/lower, no-trans/trans, unit/non-unit diag
// Factor A into LU to get well-conditioned triangles, else solve yields garbage.
// Still can give garbage if solves aren't consistent with LU factors,
// e.g., using unit diag for U, so copy lower triangle to upper triangle.
// Also used for trsm later.
lapackf77_zlacpy( "Full", &maxn, &maxn, A, &ld, LU, &ld );
lapackf77_zgetrf( &maxn, &maxn, LU, &ld, piv, &info );
for( int j = 0; j < maxn; ++j ) {
for( int i = 0; i < j; ++i ) {
*LU(i,j) = *LU(j,i);
}
}
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
magma_zsetmatrix( m, m, LU, ld, dA, ld );
magma_zsetvector( m, C, 1, dC1, 1 );
magma_zsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_ztrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZtrsv( handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
cublas_diag_const(diag[id]), m, dA, ld, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_zlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZTRSM( MagmaLeft, m, 1 ) / 1e9;
printf( "ztrsv( %c, %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]), lapacke_diag_const(diag[id]),
error, gflops/t1, gflops/t2 );
}}}
printf( "\n" );
printf( "========== Level 3 BLAS ==========\n" );
// ----- test ZGEMM
// C = alpha*A*B + beta*C, with A m*k or k*m; B k*n or n*k; C m*n
// try combinations of no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
for( int ib = 0; ib < 3; ++ib ) {
bool nta = (trans[ia] == MagmaNoTrans);
bool ntb = (trans[ib] == MagmaNoTrans);
magma_zsetmatrix( (nta ? m : k), (nta ? m : k), A, ld, dA, ld );
magma_zsetmatrix( (ntb ? k : n), (ntb ? n : k), B, ld, dB, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zgemm( trans[ia], trans[ib], m, n, k, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZgemm( handle, cublas_trans_const(trans[ia]), cublas_trans_const(trans[ib]),
m, n, k, &alpha, dA, ld, dB, ld, &beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZGEMM( m, n, k ) / 1e9;
printf( "zgemm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_trans_const(trans[ia]), lapacke_trans_const(trans[ib]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZHEMM
// C = alpha*A*B + beta*C (left) with A m*m symmetric; B,C m*n; or
// C = alpha*B*A + beta*C (right) with A n*n symmetric; B,C m*n
// try left/right, upper/lower
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
magma_zsetmatrix( m, m, A, ld, dA, ld );
magma_zsetmatrix( m, n, B, ld, dB, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zhemm( side[is], uplo[iu], m, n, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZhemm( handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
m, n, &alpha, dA, ld, dB, ld, &beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHEMM( side[is], m, n ) / 1e9;
printf( "zhemm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_side_const(side[is]), lapacke_uplo_const(uplo[iu]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZHERK
// C = alpha*A*A^H + beta*C (no-trans) with A m*k and C m*m symmetric; or
// C = alpha*A^H*A + beta*C (trans) with A k*m and C m*m symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
magma_zsetmatrix( n, k, A, ld, dA, ld );
magma_zsetmatrix( n, n, C, ld, dC1, ld );
magma_zsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zherk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZherk( handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
n, k, &dalpha, dA, ld, &dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHERK( k, n ) / 1e9;
printf( "zherk( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZHER2K
// C = alpha*A*B^H + ^alpha*B*A^H + beta*C (no-trans) with A,B n*k; C n*n symmetric; or
// C = alpha*A^H*B + ^alpha*B^H*A + beta*C (trans) with A,B k*n; C n*n symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
bool nt = (trans[it] == MagmaNoTrans);
magma_zsetmatrix( (nt ? n : k), (nt ? n : k), A, ld, dA, ld );
magma_zsetmatrix( n, n, C, ld, dC1, ld );
magma_zsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zher2k( uplo[iu], trans[it], n, k, alpha, dA, ld, dB, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZher2k( handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
n, k, &alpha, dA, ld, dB, ld, &dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHER2K( k, n ) / 1e9;
printf( "zher2k( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZTRMM
// C = alpha*A*C (left) with A m*m triangular; C m*n; or
// C = alpha*C*A (right) with A n*n triangular; C m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == MagmaLeft);
magma_zsetmatrix( (left ? m : n), (left ? m : n), A, ld, dA, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ztrmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
// note cublas does trmm out-of-place (i.e., adds output matrix C),
// but allows C=B to do in-place.
t2 = magma_sync_wtime( 0 );
cublasZtrmm( handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
cublas_trans_const(trans[it]), cublas_diag_const(diag[id]),
m, n, &alpha, dA, ld, dC2, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZTRMM( side[is], m, n ) / 1e9;
printf( "ztrmm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// ----- test ZTRSM
// solve A*X = alpha*B (left) with A m*m triangular; B m*n; or
// solve X*A = alpha*B (right) with A n*n triangular; B m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == MagmaLeft);
magma_zsetmatrix( (left ? m : n), (left ? m : n), LU, ld, dA, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ztrsm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZtrsm( handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
cublas_trans_const(trans[it]), cublas_diag_const(diag[id]),
m, n, &alpha, dA, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZTRSM( side[is], m, n ) / 1e9;
printf( "ztrsm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// cleanup
magma_free_cpu( piv );
magma_free_pinned( A );
magma_free_pinned( B );
magma_free_pinned( C );
magma_free_pinned( C2 );
magma_free_pinned( LU );
magma_free( dA );
magma_free( dB );
magma_free( dC1 );
magma_free( dC2 );
fflush( stdout );
}
if ( total_error != 0. ) {
printf( "total error %.2g -- ought to be 0 -- some test failed (see above).\n",
total_error );
}
else {
printf( "all tests passed\n" );
}
TESTING_FINALIZE();
int status = (total_error != 0.);
return status;
}
magma_int_t
magma_ssolverinfo_init( magma_s_solver_par *solver_par,
magma_s_preconditioner *precond_par ){
/*
solver_par->solver = Magma_CG;
solver_par->maxiter = 1000;
solver_par->numiter = 0;
solver_par->ortho = Magma_CGS;
solver_par->epsilon = RTOLERANCE;
solver_par->restart = 30;
solver_par->init_res = 0.;
solver_par->final_res = 0.;
solver_par->runtime = 0.;
solver_par->verbose = 0;
solver_par->info = 0;
*/
if( solver_par->verbose > 0 ){
magma_malloc_cpu( (void **)&solver_par->res_vec, sizeof(real_Double_t)
* ( (solver_par->maxiter)/(solver_par->verbose)+1) );
magma_malloc_cpu( (void **)&solver_par->timing, sizeof(real_Double_t)
*( (solver_par->maxiter)/(solver_par->verbose)+1) );
}else{
solver_par->res_vec = NULL;
solver_par->timing = NULL;
}
if( solver_par->solver == Magma_LOBPCG ){
magma_smalloc_cpu( &solver_par->eigenvalues ,
3*solver_par->num_eigenvalues );
// setup initial guess EV using lapack
// then copy to GPU
magma_int_t ev = solver_par->num_eigenvalues * solver_par->ev_length;
float *initial_guess;
magma_smalloc_cpu( &initial_guess, ev );
magma_smalloc( &solver_par->eigenvectors, ev );
magma_int_t ISEED[4] = {0,0,0,1}, ione = 1;
lapackf77_slarnv( &ione, ISEED, &ev, initial_guess );
magma_ssetmatrix( solver_par->ev_length, solver_par->num_eigenvalues,
initial_guess, solver_par->ev_length, solver_par->eigenvectors,
solver_par->ev_length );
magma_free_cpu( initial_guess );
}else{
solver_par->eigenvectors = NULL;
solver_par->eigenvalues = NULL;
}
precond_par->d.val = NULL;
precond_par->M.val = NULL;
precond_par->M.col = NULL;
precond_par->M.row = NULL;
precond_par->M.blockinfo = NULL;
precond_par->L.val = NULL;
precond_par->L.col = NULL;
precond_par->L.row = NULL;
precond_par->L.blockinfo = NULL;
precond_par->U.val = NULL;
precond_par->U.col = NULL;
precond_par->U.row = NULL;
precond_par->U.blockinfo = NULL;
precond_par->LD.val = NULL;
precond_par->LD.col = NULL;
precond_par->LD.row = NULL;
precond_par->LD.blockinfo = NULL;
precond_par->UD.val = NULL;
precond_par->UD.col = NULL;
precond_par->UD.row = NULL;
precond_par->UD.blockinfo = NULL;
return MAGMA_SUCCESS;
}
extern "C" magma_int_t
magma_zgeqrf_batched(
magma_int_t m, magma_int_t n,
magmaDoubleComplex **dA_array,
magma_int_t ldda,
magmaDoubleComplex **tau_array,
magma_int_t *info_array, magma_int_t batchCount, magma_queue_t queue)
{
#define dA(i, j) (dA + (i) + (j)*ldda) // A(i, j) means at i row, j column
magma_int_t min_mn = min(m, n);
cudaMemset(info_array, 0, batchCount*sizeof(magma_int_t));
/* Check arguments */
magma_int_t arginfo = 0;
if (m < 0)
arginfo = -1;
else if (n < 0)
arginfo = -2;
else if (ldda < max(1,m))
arginfo = -4;
if (arginfo != 0) {
magma_xerbla( __func__, -(arginfo) );
return arginfo;
}
/* Quick return if possible */
if (m == 0 || n == 0)
if(min_mn == 0 ) return arginfo;
if( m > 2048 || n > 2048 ) {
printf("=========================================================================================\n");
printf(" WARNING batched routines are designed for small sizes it might be better to use the\n Native/Hybrid classical routines if you want performance\n");
printf("=========================================================================================\n");
}
magma_int_t nb = 32;
magma_int_t nnb = 8;
magma_int_t i, k, ib=nb, jb=nnb;
magma_int_t ldw, ldt, ldr, offset;
cublasHandle_t myhandle;
cublasCreate_v2(&myhandle);
magmaDoubleComplex **dW0_displ = NULL;
magmaDoubleComplex **dW1_displ = NULL;
magmaDoubleComplex **dW2_displ = NULL;
magmaDoubleComplex **dW3_displ = NULL;
magmaDoubleComplex **dW4_displ = NULL;
magmaDoubleComplex **dW5_displ = NULL;
magmaDoubleComplex *dwork = NULL;
magmaDoubleComplex *dT = NULL;
magmaDoubleComplex *dR = NULL;
magmaDoubleComplex **dR_array = NULL;
magmaDoubleComplex **dT_array = NULL;
magmaDoubleComplex **cpuAarray = NULL;
magmaDoubleComplex **cpuTarray = NULL;
magma_malloc((void**)&dW0_displ, batchCount * sizeof(*dW0_displ));
magma_malloc((void**)&dW1_displ, batchCount * sizeof(*dW1_displ));
magma_malloc((void**)&dW2_displ, batchCount * sizeof(*dW2_displ));
magma_malloc((void**)&dW3_displ, batchCount * sizeof(*dW3_displ));
magma_malloc((void**)&dW4_displ, batchCount * sizeof(*dW4_displ)); // used in zlarfb
magma_malloc((void**)&dW5_displ, batchCount * sizeof(*dW5_displ));
magma_malloc((void**)&dR_array, batchCount * sizeof(*dR_array));
magma_malloc((void**)&dT_array, batchCount * sizeof(*dT_array));
ldt = ldr = min(nb, min_mn);
magma_zmalloc(&dwork, (2 * nb * n) * batchCount);
magma_zmalloc(&dR, ldr * n * batchCount);
magma_zmalloc(&dT, ldt * ldt * batchCount);
magma_malloc_cpu((void**) &cpuAarray, batchCount*sizeof(magmaDoubleComplex*));
magma_malloc_cpu((void**) &cpuTarray, batchCount*sizeof(magmaDoubleComplex*));
/* check allocation */
if ( dW0_displ == NULL || dW1_displ == NULL || dW2_displ == NULL || dW3_displ == NULL ||
dW4_displ == NULL || dW5_displ == NULL || dR_array == NULL || dT_array == NULL ||
dR == NULL || dT == NULL || dwork == NULL || cpuAarray == NULL || cpuTarray == NULL ) {
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dW5_displ);
magma_free(dR_array);
magma_free(dT_array);
magma_free(dR);
magma_free(dT);
magma_free(dwork);
free(cpuAarray);
free(cpuTarray);
magma_int_t info = MAGMA_ERR_DEVICE_ALLOC;
magma_xerbla( __func__, -(info) );
return info;
}
magmablas_zlaset_q(MagmaFull, ldr, n*batchCount , MAGMA_Z_ZERO, MAGMA_Z_ZERO, dR, ldr, queue);
magmablas_zlaset_q(MagmaFull, ldt, ldt*batchCount, MAGMA_Z_ZERO, MAGMA_Z_ZERO, dT, ldt, queue);
zset_pointer(dR_array, dR, 1, 0, 0, ldr*min(nb, min_mn), batchCount, queue);
zset_pointer(dT_array, dT, 1, 0, 0, ldt*min(nb, min_mn), batchCount, queue);
magma_queue_t cstream;
magmablasGetKernelStream(&cstream);
magma_int_t streamid;
const magma_int_t nbstreams=32;
magma_queue_t stream[nbstreams];
for(i=0; i<nbstreams; i++) {
magma_queue_create( &stream[i] );
}
magma_getvector( batchCount, sizeof(magmaDoubleComplex*), dA_array, 1, cpuAarray, 1);
magma_getvector( batchCount, sizeof(magmaDoubleComplex*), dT_array, 1, cpuTarray, 1);
magmablasSetKernelStream(NULL);
for(i=0; i<min_mn; i+=nb)
{
ib = min(nb, min_mn-i);
//===============================================
// panel factorization
//===============================================
magma_zdisplace_pointers(dW0_displ, dA_array, ldda, i, i, batchCount, queue);
magma_zdisplace_pointers(dW2_displ, tau_array, 1, i, 0, batchCount, queue);
//dwork is used in panel factorization and trailing matrix update
//dW4_displ, dW5_displ are used as workspace and configured inside
magma_zgeqrf_panel_batched(m-i, ib, jb,
dW0_displ, ldda,
dW2_displ,
dT_array, ldt,
dR_array, ldr,
dW1_displ,
dW3_displ,
dwork,
dW4_displ, dW5_displ,
info_array,
batchCount, myhandle, queue);
//===============================================
// end of panel
//===============================================
//direct panel matrix V in dW0_displ,
magma_zdisplace_pointers(dW0_displ, dA_array, ldda, i, i, batchCount, queue);
// copy the upper part of V into dR
zgeqrf_copy_upper_batched(ib, jb, dW0_displ, ldda, dR_array, ldr, batchCount, queue);
//===============================================
// update trailing matrix
//===============================================
//dwork is used in panel factorization and trailing matrix update
//reset dW4_displ
ldw = nb;
zset_pointer(dW4_displ, dwork, 1, 0, 0, ldw*n, batchCount, queue );
offset = ldw*n*batchCount;
zset_pointer(dW5_displ, dwork + offset, 1, 0, 0, ldw*n, batchCount, queue );
if( (n-ib-i) > 0)
{
// set the diagonal of v as one and the upper triangular part as zero
magmablas_zlaset_batched(MagmaUpper, ib, ib, MAGMA_Z_ZERO, MAGMA_Z_ONE, dW0_displ, ldda, batchCount, queue);
magma_zdisplace_pointers(dW2_displ, tau_array, 1, i, 0, batchCount, queue);
// it is faster since it is using BLAS-3 GEMM routines, different from lapack implementation
magma_zlarft_batched(m-i, ib, 0,
dW0_displ, ldda,
dW2_displ,
dT_array, ldt,
dW4_displ, nb*ldt,
batchCount, myhandle, queue);
// perform C = (I-V T^H V^H) * C, C is the trailing matrix
//-------------------------------------------
// USE STREAM GEMM
//-------------------------------------------
if( (m-i) > 100 && (n-i-ib) > 100)
{
// But since the code use the NULL stream everywhere,
// so I don't need it, because the NULL stream do the sync by itself
//magma_device_sync();
for(k=0; k<batchCount; k++)
{
streamid = k%nbstreams;
magmablasSetKernelStream(stream[streamid]);
// the stream gemm must take cpu pointer
magma_zlarfb_gpu_gemm(MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-i, n-i-ib, ib,
cpuAarray[k] + i + i * ldda, ldda,
cpuTarray[k], ldt,
cpuAarray[k] + i + (i+ib) * ldda, ldda,
dwork + nb * n * k, -1,
dwork + nb * n * batchCount + nb * n * k, -1);
}
// need to synchronise to be sure that panel does not start before
// finishing the update at least of the next panel
// BUT no need for it as soon as the other portion of the code
// use the NULL stream which do the sync by itself
//magma_device_sync();
magmablasSetKernelStream(NULL);
}
//-------------------------------------------
// USE BATCHED GEMM
//-------------------------------------------
else
{
//direct trailing matrix in dW1_displ
magma_zdisplace_pointers(dW1_displ, dA_array, ldda, i, i+ib, batchCount, queue);
magma_zlarfb_gemm_batched(
MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-i, n-i-ib, ib,
(const magmaDoubleComplex**)dW0_displ, ldda,
(const magmaDoubleComplex**)dT_array, ldt,
dW1_displ, ldda,
dW4_displ, ldw,
dW5_displ, ldw,
batchCount, myhandle, queue);
}
}// update the trailing matrix
//===============================================
// copy dR back to V after the trailing matrix update
magmablas_zlacpy_batched(MagmaUpper, ib, ib, dR_array, ldr, dW0_displ, ldda, batchCount, queue);
}
for(k=0; k<nbstreams; k++) {
magma_queue_destroy( stream[k] );
}
magmablasSetKernelStream(cstream);
cublasDestroy_v2(myhandle);
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dW5_displ);
magma_free(dR_array);
magma_free(dT_array);
magma_free(dR);
magma_free(dT);
magma_free(dwork);
free(cpuAarray);
free(cpuTarray);
return arginfo;
}
extern "C" magma_int_t magma_ssytrd_sb2st(magma_int_t threads, char uplo, magma_int_t n, magma_int_t nb, magma_int_t Vblksiz,
float *A, magma_int_t lda, float *D, float *E,
float *V, magma_int_t ldv, float *TAU, magma_int_t compT, float *T, magma_int_t ldt)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
Arguments
=========
THREADS (input) INTEGER
Specifies the number of pthreads used.
THREADS > 0
UPLO (input) CHARACTER*1
= 'U': Upper triangles of A is stored;
= 'L': Lower triangles of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
NB (input) INTEGER
The order of the band matrix A. N >= NB >= 0.
VBLKSIZ (input) INTEGER
The size of the block of householder vectors applied at once.
A (input/workspace) REAL array, dimension (LDA, N)
On entry the band matrix stored in the following way:
LDA (input) INTEGER
The leading dimension of the array A. LDA >= 2*NB.
D (output) DOUBLE array, dimension (N)
The diagonal elements of the tridiagonal matrix T:
D(i) = A(i,i).
E (output) DOUBLE array, dimension (N-1)
The off-diagonal elements of the tridiagonal matrix T:
E(i) = A(i,i+1) if UPLO = 'U', E(i) = A(i+1,i) if UPLO = 'L'.
V (output) REAL array, dimension (BLKCNT, LDV, VBLKSIZ)
On exit it contains the blocks of householder reflectors
BLKCNT is the number of block and it is returned by the funtion MAGMA_BULGE_GET_BLKCNT.
LDV (input) INTEGER
The leading dimension of V.
LDV > NB + VBLKSIZ + 1
TAU (output) REAL dimension(BLKCNT, VBLKSIZ)
???
COMPT (input) INTEGER
if COMPT = 0 T is not computed
if COMPT = 1 T is computed
T (output) REAL dimension(LDT *)
if COMPT = 1 on exit contains the matrices T needed for Q2
if COMPT = 0 T is not referenced
LDT (input) INTEGER
The leading dimension of T.
LDT > Vblksiz
INFO (output) INTEGER ????????????????????????????????????????????????????????????????????????????????????
= 0: successful exit
===================================================================== */
#ifdef ENABLE_TIMER
real_Double_t timeblg=0.0;
#endif
//char uplo_[2] = {uplo, 0};
magma_int_t mklth = threads;
magma_int_t INgrsiz=1;
magma_int_t blkcnt = magma_bulge_get_blkcnt(n, nb, Vblksiz);
magma_int_t nbtiles = magma_ceildiv(n, nb);
memset(T, 0, blkcnt*ldt*Vblksiz*sizeof(float));
memset(TAU, 0, blkcnt*Vblksiz*sizeof(float));
memset(V, 0, blkcnt*ldv*Vblksiz*sizeof(float));
magma_int_t* prog;
magma_malloc_cpu((void**) &prog, (2*nbtiles+threads+10)*sizeof(magma_int_t));
memset(prog, 0, (2*nbtiles+threads+10)*sizeof(magma_int_t));
magma_sbulge_id_data* arg;
magma_malloc_cpu((void**) &arg, threads*sizeof(magma_sbulge_id_data));
pthread_t* thread_id;
magma_malloc_cpu((void**) &thread_id, threads*sizeof(pthread_t));
pthread_attr_t thread_attr;
magma_setlapack_numthreads(1);
magma_sbulge_data data_bulge(threads, n, nb, nbtiles, INgrsiz, Vblksiz, compT,
A, lda, V, ldv, TAU, T, ldt, prog);
// Set one thread per core
pthread_attr_init(&thread_attr);
pthread_attr_setscope(&thread_attr, PTHREAD_SCOPE_SYSTEM);
pthread_setconcurrency(threads);
//timing
#ifdef ENABLE_TIMER
timeblg = magma_wtime();
#endif
// Launch threads
for (magma_int_t thread = 1; thread < threads; thread++)
{
arg[thread] = magma_sbulge_id_data(thread, &data_bulge);
pthread_create(&thread_id[thread], &thread_attr, magma_ssytrd_sb2st_parallel_section, &arg[thread]);
}
arg[0] = magma_sbulge_id_data(0, &data_bulge);
magma_ssytrd_sb2st_parallel_section(&arg[0]);
// Wait for completion
for (magma_int_t thread = 1; thread < threads; thread++)
{
void *exitcodep;
pthread_join(thread_id[thread], &exitcodep);
}
// timing
#ifdef ENABLE_TIMER
timeblg = magma_wtime()-timeblg;
printf(" time BULGE+T = %f \n" ,timeblg);
#endif
magma_free_cpu(thread_id);
magma_free_cpu(arg);
magma_free_cpu(prog);
magma_setlapack_numthreads(mklth);
/*================================================
* store resulting diag and lower diag D and E
* note that D and E are always real
*================================================*/
/* Make diagonal and superdiagonal elements real,
* storing them in D and E
*/
/* In real case, the off diagonal element are
* not necessary real. we have to make off-diagonal
* elements real and copy them to E.
* When using HouseHolder elimination,
* the SLARFG give us a real as output so, all the
* diagonal/off-diagonal element except the last one are already
* real and thus we need only to take the abs of the last
* one.
* */
#if defined(PRECISION_z) || defined(PRECISION_c)
if(uplo==MagmaLower){
for (magma_int_t i=0; i < n-1 ; i++)
{
D[i] = MAGMA_S_REAL(A[i*lda ]);
E[i] = MAGMA_S_REAL(A[i*lda+1]);
}
D[n-1] = MAGMA_S_REAL(A[(n-1)*lda]);
} else { /* MagmaUpper not tested yet */
for (magma_int_t i=0; i<n-1; i++)
{
D[i] = MAGMA_S_REAL(A[i*lda+nb]);
E[i] = MAGMA_S_REAL(A[i*lda+nb-1]);
}
D[n-1] = MAGMA_S_REAL(A[(n-1)*lda+nb]);
} /* end MagmaUpper */
#else
if( uplo == MagmaLower ){
for (magma_int_t i=0; i < n-1; i++) {
D[i] = A[i*lda]; // diag
E[i] = A[i*lda+1]; //lower diag
}
D[n-1] = A[(n-1)*lda];
} else {
for (magma_int_t i=0; i < n-1; i++) {
D[i] = A[i*lda+nb]; // diag
E[i] = A[i*lda+nb-1]; //lower diag
}
D[n-1] = A[(n-1)*lda+nb];
}
#endif
return MAGMA_SUCCESS;
}
extern "C" magma_int_t
magma_dsolverinfo_init(
magma_d_solver_par *solver_par,
magma_d_preconditioner *precond_par,
magma_queue_t queue )
{
magma_int_t info = 0;
solver_par->res_vec = NULL;
solver_par->timing = NULL;
solver_par->eigenvectors = NULL;
solver_par->eigenvalues = NULL;
if( solver_par->maxiter == 0 )
solver_par->maxiter = 1000;
if( solver_par->version == 0 )
solver_par->version = 0;
if( solver_par->restart == 0 )
solver_par->restart = 30;
if( solver_par->solver == 0 )
solver_par->solver = Magma_CG;
if ( solver_par->verbose > 0 ) {
CHECK( magma_malloc_cpu( (void **)&solver_par->res_vec, sizeof(real_Double_t)
* ( (solver_par->maxiter)/(solver_par->verbose)+1) ));
CHECK( magma_malloc_cpu( (void **)&solver_par->timing, sizeof(real_Double_t)
*( (solver_par->maxiter)/(solver_par->verbose)+1) ));
} else {
solver_par->res_vec = NULL;
solver_par->timing = NULL;
}
precond_par->d.val = NULL;
precond_par->d2.val = NULL;
precond_par->work1.val = NULL;
precond_par->work2.val = NULL;
precond_par->M.val = NULL;
precond_par->M.col = NULL;
precond_par->M.row = NULL;
precond_par->M.blockinfo = NULL;
precond_par->L.val = NULL;
precond_par->L.col = NULL;
precond_par->L.row = NULL;
precond_par->L.blockinfo = NULL;
precond_par->U.val = NULL;
precond_par->U.col = NULL;
precond_par->U.row = NULL;
precond_par->U.blockinfo = NULL;
precond_par->LD.val = NULL;
precond_par->LD.col = NULL;
precond_par->LD.row = NULL;
precond_par->LD.blockinfo = NULL;
precond_par->UD.val = NULL;
precond_par->UD.col = NULL;
precond_par->UD.row = NULL;
precond_par->UD.blockinfo = NULL;
precond_par->cuinfoL = NULL;
precond_par->cuinfoU = NULL;
cleanup:
if( info != 0 ){
magma_free( solver_par->timing );
magma_free( solver_par->res_vec );
}
return info;
}
/***************************************************************************//**
Purpose
-------
DGEQRF computes a QR factorization of a real M-by-N matrix A:
A = Q * R.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
dA_array Array of pointers, dimension (batchCount).
Each is a DOUBLE PRECISION array on the GPU, dimension (LDDA,N)
On entry, the M-by-N matrix A.
On exit, the elements on and above the diagonal of the array
contain the min(M,N)-by-N upper trapezoidal matrix R (R is
upper triangular if m >= n); the elements below the diagonal,
with the array TAU, represent the orthogonal matrix Q as a
product of min(m,n) elementary reflectors (see Further
Details).
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,M).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param[in,out]
dR_array Array of pointers, dimension (batchCount).
Each is a DOUBLE PRECISION array on the GPU, dimension (LDDR, N/NB)
dR should be of size (LDDR, N) when provide_RT > 0 and
of size (LDDT, NB) otherwise. NB is the local blocking size.
On exit, the elements of R are stored in dR only when provide_RT > 0.
@param[in]
lddr INTEGER
The leading dimension of the array dR.
LDDR >= min(M,N) when provide_RT == 1
otherwise LDDR >= min(NB, min(M,N)).
NB is the local blocking size.
To benefit from coalescent memory accesses LDDR must be
divisible by 16.
@param[in,out]
dT_array Array of pointers, dimension (batchCount).
Each is a DOUBLE PRECISION array on the GPU, dimension (LDDT, N/NB)
dT should be of size (LDDT, N) when provide_RT > 0 and
of size (LDDT, NB) otherwise. NB is the local blocking size.
On exit, the elements of T are stored in dT only when provide_RT > 0.
@param[in]
lddt INTEGER
The leading dimension of the array dT.
LDDT >= min(NB,min(M,N)). NB is the local blocking size.
To benefit from coalescent memory accesses LDDR must be
divisible by 16.
@param[out]
dtau_array Array of pointers, dimension (batchCount).
Each is a DOUBLE PRECISION array, dimension (min(M,N))
The scalar factors of the elementary reflectors (see Further
Details).
@param[in]
provide_RT INTEGER
provide_RT = 0 no R and no T in output.
dR and dT are used as local workspace to store the R and T of each step.
provide_RT = 1 the whole R of size (min(M,N), N) and the nbxnb block of T are provided in output.
provide_RT = 2 the nbxnb diag block of R and of T are provided in output.
@param[out]
info_array Array of INTEGERs, dimension (batchCount), for corresponding matrices.
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
@param[in]
batchCount INTEGER
The number of matrices to operate on.
@param[in]
queue magma_queue_t
Queue to execute in.
Further Details
---------------
The matrix Q is represented as a product of elementary reflectors
Q = H(1) H(2) . . . H(k), where k = min(m,n).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i-1) = 0 and v(i) = 1; v(i+1:m) is stored on exit in A(i+1:m,i),
and tau in TAU(i).
@ingroup magma_geqrf_batched
*******************************************************************************/
extern "C" magma_int_t
magma_dgeqrf_expert_batched(
magma_int_t m, magma_int_t n,
double **dA_array, magma_int_t ldda,
double **dR_array, magma_int_t lddr,
double **dT_array, magma_int_t lddt,
double **dtau_array, magma_int_t provide_RT,
magma_int_t *info_array, magma_int_t batchCount, magma_queue_t queue)
{
#define dA(i, j) (dA + (i) + (j)*ldda)
/* Local Parameter */
magma_int_t nb = magma_get_dgeqrf_batched_nb(m);
magma_int_t nnb = 8;
magma_int_t min_mn = min(m, n);
/* Check arguments */
cudaMemset(info_array, 0, batchCount*sizeof(magma_int_t));
magma_int_t arginfo = 0;
if (m < 0)
arginfo = -1;
else if (n < 0)
arginfo = -2;
else if (ldda < max(1,m))
arginfo = -4;
else if (lddr < min_mn && provide_RT == 1)
arginfo = -6;
else if (lddr < min(min_mn, nb))
arginfo = -6;
else if (lddt < min(min_mn, nb))
arginfo = -8;
if (arginfo != 0) {
magma_xerbla( __func__, -(arginfo) );
return arginfo;
}
/* Quick return if possible */
if (m == 0 || n == 0)
if (min_mn == 0 ) return arginfo;
if ( m > 2048 || n > 2048 ) {
printf("=========================================================================================\n");
printf(" WARNING batched routines are designed for small sizes it might be better to use the\n Native/Hybrid classical routines if you want performance\n");
printf("=========================================================================================\n");
}
magma_int_t i, k, ib=nb, jb=nnb, offset_RT=0, use_stream;
magma_int_t ldw, offset;
double **dW0_displ = NULL;
double **dW1_displ = NULL;
double **dW2_displ = NULL;
double **dW3_displ = NULL;
double **dW4_displ = NULL;
double **dW5_displ = NULL;
double **dR_displ = NULL;
double **dT_displ = NULL;
double *dwork = NULL;
double **cpuAarray = NULL;
double **cpuTarray = NULL;
magma_malloc((void**)&dW0_displ, batchCount * sizeof(*dW0_displ));
magma_malloc((void**)&dW1_displ, batchCount * sizeof(*dW1_displ));
magma_malloc((void**)&dW2_displ, batchCount * sizeof(*dW2_displ));
magma_malloc((void**)&dW3_displ, batchCount * sizeof(*dW3_displ));
magma_malloc((void**)&dW4_displ, batchCount * sizeof(*dW4_displ));
magma_malloc((void**)&dW5_displ, batchCount * sizeof(*dW5_displ));
magma_malloc((void**)&dR_displ, batchCount * sizeof(*dR_displ));
magma_malloc((void**)&dT_displ, batchCount * sizeof(*dT_displ));
magma_dmalloc(&dwork, (2 * nb * n) * batchCount);
magma_malloc_cpu((void**) &cpuAarray, batchCount*sizeof(double*));
magma_malloc_cpu((void**) &cpuTarray, batchCount*sizeof(double*));
/* check allocation */
if ( dW0_displ == NULL || dW1_displ == NULL || dW2_displ == NULL ||
dW3_displ == NULL || dW4_displ == NULL || dW5_displ == NULL ||
dR_displ == NULL || dT_displ == NULL || dwork == NULL ||
cpuAarray == NULL || cpuTarray == NULL ) {
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dW5_displ);
magma_free(dR_displ);
magma_free(dT_displ);
magma_free(dwork);
magma_free_cpu(cpuAarray);
magma_free_cpu(cpuTarray);
magma_int_t info = MAGMA_ERR_DEVICE_ALLOC;
magma_xerbla( __func__, -(info) );
return info;
}
magma_ddisplace_pointers(dR_displ, dR_array, lddr, 0, 0, batchCount, queue);
magma_ddisplace_pointers(dT_displ, dT_array, lddt, 0, 0, batchCount, queue);
// set dwork to zero because our GEMM routine does propagate NAN when C=betaC+alphaA*B and beta=0
magmablas_dlaset_q( MagmaFull, 2*nb, n*batchCount, MAGMA_D_ZERO, MAGMA_D_ZERO, dwork, 2*nb, queue );
// set dR and dT to zero. if provide_RT == 0 only a tile of size nbxnb is used and overwritten at each step
magmablas_dlaset_batched( MagmaFull, lddr, (provide_RT > 0 ? n:min(min_mn,nb)), MAGMA_D_ZERO, MAGMA_D_ZERO, dR_displ, lddr, batchCount, queue );
magmablas_dlaset_batched( MagmaFull, lddt, (provide_RT > 0 ? n:min(min_mn,nb)), MAGMA_D_ZERO, MAGMA_D_ZERO, dT_displ, lddt, batchCount, queue );
/*
if ( provide_RT > 0 )
{
magmablas_dlaset_q( MagmaFull, lddr, n*batchCount, MAGMA_D_ZERO, MAGMA_D_ZERO, dR, lddr, queue );
magmablas_dlaset_q( MagmaFull, lddt, n*batchCount, MAGMA_D_ZERO, MAGMA_D_ZERO, dT, lddt, queue );
}
else
{
magmablas_dlaset_q( MagmaFull, lddr, nb*batchCount, MAGMA_D_ZERO, MAGMA_D_ZERO, dR, lddr, queue );
magmablas_dlaset_q( MagmaFull, lddt, nb*batchCount, MAGMA_D_ZERO, MAGMA_D_ZERO, dT, lddt, queue );
}
*/
magma_int_t streamid;
const magma_int_t nbstreams=10;
magma_queue_t queues[nbstreams];
for (i=0; i < nbstreams; i++) {
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[i] );
}
magma_getvector( batchCount, sizeof(double*), dA_array, 1, cpuAarray, 1, queue);
magma_getvector( batchCount, sizeof(double*), dT_array, 1, cpuTarray, 1, queue);
for (i=0; i < min_mn; i += nb)
{
ib = min(nb, min_mn-i);
//===============================================
// panel factorization
//===============================================
magma_ddisplace_pointers(dW0_displ, dA_array, ldda, i, i, batchCount, queue);
magma_ddisplace_pointers(dW2_displ, dtau_array, 1, i, 0, batchCount, queue);
if ( provide_RT > 0 )
{
offset_RT = i;
magma_ddisplace_pointers(dR_displ, dR_array, lddr, (provide_RT == 1 ? offset_RT:0), offset_RT, batchCount, queue);
magma_ddisplace_pointers(dT_displ, dT_array, lddt, 0, offset_RT, batchCount, queue);
}
//dwork is used in panel factorization and trailing matrix update
//dW4_displ, dW5_displ are used as workspace and configured inside
magma_dgeqrf_panel_batched(m-i, ib, jb,
dW0_displ, ldda,
dW2_displ,
dT_displ, lddt,
dR_displ, lddr,
dW1_displ,
dW3_displ,
dwork,
dW4_displ, dW5_displ,
info_array,
batchCount, queue);
//===============================================
// end of panel
//===============================================
//===============================================
// update trailing matrix
//===============================================
if ( (n-ib-i) > 0)
{
//dwork is used in panel factorization and trailing matrix update
//reset dW4_displ
ldw = nb;
magma_dset_pointer( dW4_displ, dwork, 1, 0, 0, ldw*n, batchCount, queue );
offset = ldw*n*batchCount;
magma_dset_pointer( dW5_displ, dwork + offset, 1, 0, 0, ldw*n, batchCount, queue );
// set the diagonal of v as one and the upper triangular part as zero already set inside geqrf_panel
//magmablas_dlaset_batched( MagmaUpper, ib, ib, MAGMA_D_ZERO, MAGMA_D_ONE, dW0_displ, ldda, batchCount, queue );
//magma_ddisplace_pointers(dW2_displ, dtau_array, 1, i, 0, batchCount, queue);
// it is faster since it is using BLAS-3 GEMM routines, different from lapack implementation
magma_dlarft_batched(m-i, ib, 0,
dW0_displ, ldda,
dW2_displ,
dT_displ, lddt,
dW4_displ, nb*lddt,
batchCount, queue);
// perform C = (I-V T^H V^H) * C, C is the trailing matrix
//-------------------------------------------
// USE STREAM GEMM
//-------------------------------------------
use_stream = magma_drecommend_cublas_gemm_stream(MagmaNoTrans, MagmaNoTrans, m-i-ib, n-i-ib, ib);
if ( use_stream )
{
magma_queue_sync(queue);
for (k=0; k < batchCount; k++)
{
streamid = k%nbstreams;
// the queue gemm must take cpu pointer
magma_dlarfb_gpu_gemm( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-i, n-i-ib, ib,
cpuAarray[k] + i + i * ldda, ldda,
cpuTarray[k] + offset_RT*lddt, lddt,
cpuAarray[k] + i + (i+ib) * ldda, ldda,
dwork + nb * n * k, -1,
dwork + nb * n * batchCount + nb * n * k, -1, queues[streamid] );
}
// need to synchronise to be sure that panel does not start before
// finishing the update at least of the next panel
// if queue is NULL, no need to sync
if ( queue != NULL ) {
for (magma_int_t s=0; s < nbstreams; s++)
magma_queue_sync(queues[s]);
}
}
//-------------------------------------------
// USE BATCHED GEMM
//-------------------------------------------
else
{
//direct trailing matrix in dW1_displ
magma_ddisplace_pointers(dW1_displ, dA_array, ldda, i, i+ib, batchCount, queue);
magma_dlarfb_gemm_batched(
MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-i, n-i-ib, ib,
(const double**)dW0_displ, ldda,
(const double**)dT_displ, lddt,
dW1_displ, ldda,
dW4_displ, ldw,
dW5_displ, ldw,
batchCount, queue );
}
}// update the trailing matrix
//===============================================
// copy dR back to V after the trailing matrix update,
// only when provide_RT=0 otherwise the nbxnb block of V is set to diag=1/0
// The upper portion of V could be set totaly to 0 here
if ( provide_RT == 0 )
{
magmablas_dlacpy_batched( MagmaUpper, ib, ib, dR_displ, lddr, dW0_displ, ldda, batchCount, queue );
}
}
magma_queue_sync(queue);
for (k=0; k < nbstreams; k++) {
magma_queue_destroy( queues[k] );
}
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dW5_displ);
magma_free(dR_displ);
magma_free(dT_displ);
magma_free(dwork);
magma_free_cpu(cpuAarray);
magma_free_cpu(cpuTarray);
return arginfo;
}
/**
Purpose
-------
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangles of A is stored;
- = MagmaLower: Lower triangles of A is stored.
@param[in]
n INTEGER
The order of the matrix A. N >= 0.
@param[in]
nb INTEGER
The order of the band matrix A. N >= NB >= 0.
@param[in]
Vblksiz INTEGER
The size of the block of householder vectors applied at once.
@param[in]
A (workspace) COMPLEX_16 array, dimension (LDA, N)
On entry the band matrix stored in the following way:
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= 2*NB.
@param[out]
d DOUBLE array, dimension (N)
The diagonal elements of the tridiagonal matrix T:
D(i) = A(i,i).
@param[out]
e DOUBLE array, dimension (N-1)
The off-diagonal elements of the tridiagonal matrix T:
E(i) = A(i,i+1) if UPLO = MagmaUpper, E(i) = A(i+1,i) if UPLO = MagmaLower.
@param[out]
V COMPLEX_16 array, dimension (BLKCNT, LDV, VBLKSIZ)
On exit it contains the blocks of householder reflectors
BLKCNT is the number of block and it is returned by the funtion MAGMA_BULGE_GET_BLKCNT.
@param[in]
ldv INTEGER
The leading dimension of V.
LDV > NB + VBLKSIZ + 1
@param[out]
TAU COMPLEX_16 dimension(BLKCNT, VBLKSIZ)
???
@param[in]
compT INTEGER
if COMPT = 0 T is not computed
if COMPT = 1 T is computed
@param[out]
T COMPLEX_16 dimension(LDT *)
if COMPT = 1 on exit contains the matrices T needed for Q2
if COMPT = 0 T is not referenced
@param[in]
ldt INTEGER
The leading dimension of T.
LDT > Vblksiz
@ingroup magma_zheev_2stage
********************************************************************/
extern "C" magma_int_t
magma_zhetrd_hb2st(
magma_uplo_t uplo, magma_int_t n, magma_int_t nb, magma_int_t Vblksiz,
magmaDoubleComplex *A, magma_int_t lda, double *d, double *e,
magmaDoubleComplex *V, magma_int_t ldv, magmaDoubleComplex *TAU,
magma_int_t compT, magmaDoubleComplex *T, magma_int_t ldt)
{
#ifdef ENABLE_TIMER
real_Double_t timeblg=0.0;
#endif
magma_int_t threads = magma_get_parallel_numthreads();
magma_int_t mklth = magma_get_lapack_numthreads();
magma_set_lapack_numthreads(1);
//const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t INgrsiz=1;
magma_int_t blkcnt = magma_bulge_get_blkcnt(n, nb, Vblksiz);
magma_int_t nbtiles = magma_ceildiv(n, nb);
memset(T, 0, blkcnt*ldt*Vblksiz*sizeof(magmaDoubleComplex));
memset(TAU, 0, blkcnt*Vblksiz*sizeof(magmaDoubleComplex));
memset(V, 0, blkcnt*ldv*Vblksiz*sizeof(magmaDoubleComplex));
magma_int_t* prog;
magma_malloc_cpu((void**) &prog, (2*nbtiles+threads+10)*sizeof(magma_int_t));
memset(prog, 0, (2*nbtiles+threads+10)*sizeof(magma_int_t));
magma_zbulge_id_data* arg;
magma_malloc_cpu((void**) &arg, threads*sizeof(magma_zbulge_id_data));
pthread_t* thread_id;
magma_malloc_cpu((void**) &thread_id, threads*sizeof(pthread_t));
pthread_attr_t thread_attr;
magma_zbulge_data data_bulge(threads, n, nb, nbtiles, INgrsiz, Vblksiz, compT,
A, lda, V, ldv, TAU, T, ldt, prog);
// Set one thread per core
pthread_attr_init(&thread_attr);
pthread_attr_setscope(&thread_attr, PTHREAD_SCOPE_SYSTEM);
pthread_setconcurrency(threads);
//timing
#ifdef ENABLE_TIMER
timeblg = magma_wtime();
#endif
// Launch threads
for (magma_int_t thread = 1; thread < threads; thread++) {
arg[thread] = magma_zbulge_id_data(thread, &data_bulge);
pthread_create(&thread_id[thread], &thread_attr, magma_zhetrd_hb2st_parallel_section, &arg[thread]);
}
arg[0] = magma_zbulge_id_data(0, &data_bulge);
magma_zhetrd_hb2st_parallel_section(&arg[0]);
// Wait for completion
for (magma_int_t thread = 1; thread < threads; thread++) {
void *exitcodep;
pthread_join(thread_id[thread], &exitcodep);
}
// timing
#ifdef ENABLE_TIMER
timeblg = magma_wtime()-timeblg;
printf(" time BULGE+T = %f\n", timeblg);
#endif
magma_free_cpu(thread_id);
magma_free_cpu(arg);
magma_free_cpu(prog);
magma_set_lapack_numthreads(mklth);
/*================================================
* store resulting diag and lower diag d and e
* note that d and e are always real
*================================================*/
/* Make diagonal and superdiagonal elements real,
* storing them in d and e
*/
/* In complex case, the off diagonal element are
* not necessary real. we have to make off-diagonal
* elements real and copy them to e.
* When using HouseHolder elimination,
* the ZLARFG give us a real as output so, all the
* diagonal/off-diagonal element except the last one are already
* real and thus we need only to take the abs of the last
* one.
* */
#if defined(PRECISION_z) || defined(PRECISION_c)
if (uplo == MagmaLower) {
for (magma_int_t i=0; i < n-1; i++) {
d[i] = MAGMA_Z_REAL( A[i*lda ] );
e[i] = MAGMA_Z_REAL( A[i*lda+1] );
}
d[n-1] = MAGMA_Z_REAL(A[(n-1)*lda]);
} else { /* MagmaUpper not tested yet */
for (magma_int_t i=0; i < n-1; i++) {
d[i] = MAGMA_Z_REAL( A[i*lda+nb] );
e[i] = MAGMA_Z_REAL( A[i*lda+nb-1] );
}
d[n-1] = MAGMA_Z_REAL(A[(n-1)*lda+nb]);
} /* end MagmaUpper */
#else
if ( uplo == MagmaLower ) {
for (magma_int_t i=0; i < n-1; i++) {
d[i] = A[i*lda]; // diag
e[i] = A[i*lda+1]; // lower diag
}
d[n-1] = A[(n-1)*lda];
} else {
for (magma_int_t i=0; i < n-1; i++) {
d[i] = A[i*lda+nb]; // diag
e[i] = A[i*lda+nb-1]; // lower diag
}
d[n-1] = A[(n-1)*lda+nb];
}
#endif
return MAGMA_SUCCESS;
}
/**
Purpose
-------
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangles of A is stored;
- = MagmaLower: Lower triangles of A is stored.
@param[in]
n INTEGER
The order of the matrix A. n >= 0.
@param[in]
nb INTEGER
The order of the band matrix A. n >= nb >= 0.
@param[in]
Vblksiz INTEGER
The size of the block of householder vectors applied at once.
@param[in]
A (workspace) DOUBLE PRECISION array, dimension (lda, n)
On entry the band matrix stored in the following way:
@param[in]
lda INTEGER
The leading dimension of the array A. lda >= 2*nb.
@param[out]
d DOUBLE array, dimension (n)
The diagonal elements of the tridiagonal matrix T:
D(i) = A(i,i).
@param[out]
e DOUBLE array, dimension (n-1)
The off-diagonal elements of the tridiagonal matrix T:
E(i) = A(i,i+1) if UPLO = MagmaUpper, E(i) = A(i+1,i) if UPLO = MagmaLower.
@param[out]
V DOUBLE PRECISION array, dimension (BLKCNT, LDV, VBLKSIZ)
On exit it contains the blocks of householder reflectors
BLKCNT is the number of block and it is returned by the funtion MAGMA_BULGE_GET_BLKCNT.
@param[in]
ldv INTEGER
The leading dimension of V.
LDV > nb + VBLKSIZ + 1
@param[out]
TAU DOUBLE PRECISION dimension(BLKCNT, VBLKSIZ)
???
@param[in]
wantz INTEGER
if COMPT = 0 T is not computed
if COMPT = 1 T is computed
@param[out]
T DOUBLE PRECISION dimension(LDT *)
if COMPT = 1 on exit contains the matrices T needed for Q2
if COMPT = 0 T is not referenced
@param[in]
ldt INTEGER
The leading dimension of T.
LDT > Vblksiz
@ingroup magma_dsyev_2stage
********************************************************************/
extern "C" magma_int_t
magma_dsytrd_sb2st(
magma_uplo_t uplo, magma_int_t n, magma_int_t nb, magma_int_t Vblksiz,
double *A, magma_int_t lda, double *d, double *e,
double *V, magma_int_t ldv, double *TAU,
magma_int_t wantz, double *T, magma_int_t ldt)
{
#ifdef ENABLE_TIMER
real_Double_t timeblg=0.0;
#endif
magma_int_t parallel_threads = magma_get_parallel_numthreads();
magma_int_t mklth = magma_get_lapack_numthreads();
magma_int_t ompth = magma_get_omp_numthreads();
//magma_set_omp_numthreads(1);
//magma_set_lapack_numthreads(1);
magma_int_t blkcnt, sizTAU2, sizT2, sizV2;
magma_dbulge_getstg2size(n, nb, wantz,
Vblksiz, ldv, ldt, &blkcnt,
&sizTAU2, &sizT2, &sizV2);
memset(T, 0, sizT2*sizeof(double));
memset(TAU, 0, sizTAU2*sizeof(double));
memset(V, 0, sizV2*sizeof(double));
magma_int_t INgrsiz=1;
magma_int_t nbtiles = magma_ceildiv(n, nb);
volatile magma_int_t* prog;
magma_malloc_cpu((void**) &prog, (2*nbtiles+parallel_threads+10)*sizeof(magma_int_t));
memset((void *) prog, 0, (2*nbtiles+parallel_threads+10)*sizeof(magma_int_t));
magma_dbulge_id_data* arg;
magma_malloc_cpu((void**) &arg, parallel_threads*sizeof(magma_dbulge_id_data));
pthread_t* thread_id;
magma_malloc_cpu((void**) &thread_id, parallel_threads*sizeof(pthread_t));
pthread_attr_t thread_attr;
magma_dbulge_data data_bulge;
magma_dbulge_data_init(&data_bulge, parallel_threads, n, nb, nbtiles, INgrsiz, Vblksiz, wantz,
A, lda, V, ldv, TAU, T, ldt, prog);
// Set one thread per core
pthread_attr_init(&thread_attr);
pthread_attr_setscope(&thread_attr, PTHREAD_SCOPE_SYSTEM);
pthread_setconcurrency(parallel_threads);
//timing
#ifdef ENABLE_TIMER
timeblg = magma_wtime();
#endif
// Launch threads
for (magma_int_t thread = 1; thread < parallel_threads; thread++) {
magma_dbulge_id_data_init(&(arg[thread]), thread, &data_bulge);
pthread_create(&thread_id[thread], &thread_attr, magma_dsytrd_sb2st_parallel_section, &arg[thread]);
}
magma_dbulge_id_data_init(&(arg[0]), 0, &data_bulge);
magma_dsytrd_sb2st_parallel_section(&arg[0]);
// Wait for completion
for (magma_int_t thread = 1; thread < parallel_threads; thread++) {
void *exitcodep;
pthread_join(thread_id[thread], &exitcodep);
}
// timing
#ifdef ENABLE_TIMER
timeblg = magma_wtime()-timeblg;
printf(" time BULGE+T = %f\n", timeblg);
#endif
magma_free_cpu(thread_id);
magma_free_cpu(arg);
magma_free_cpu((void *) prog);
magma_dbulge_data_destroy(&data_bulge);
magma_set_omp_numthreads(ompth);
magma_set_lapack_numthreads(mklth);
/*================================================
* store resulting diag and lower diag d and e
* note that d and e are always real
*================================================*/
/* Make diagonal and superdiagonal elements real,
* storing them in d and e
*/
/* In real case, the off diagonal element are
* not necessary real. we have to make off-diagonal
* elements real and copy them to e.
* When using HouseHolder elimination,
* the DLARFG give us a real as output so, all the
* diagonal/off-diagonal element except the last one are already
* real and thus we need only to take the abs of the last
* one.
* */
#ifdef COMPLEX
if (uplo == MagmaLower) {
for (magma_int_t i=0; i < n-1; i++) {
d[i] = MAGMA_D_REAL( A[i*lda ] );
e[i] = MAGMA_D_REAL( A[i*lda+1] );
}
d[n-1] = MAGMA_D_REAL(A[(n-1)*lda]);
} else { /* MagmaUpper not tested yet */
for (magma_int_t i=0; i < n-1; i++) {
d[i] = MAGMA_D_REAL( A[i*lda+nb] );
e[i] = MAGMA_D_REAL( A[i*lda+nb-1] );
}
d[n-1] = MAGMA_D_REAL(A[(n-1)*lda+nb]);
} /* end MagmaUpper */
#else
if ( uplo == MagmaLower ) {
for (magma_int_t i=0; i < n-1; i++) {
d[i] = A[i*lda]; // diag
e[i] = A[i*lda+1]; // lower diag
}
d[n-1] = A[(n-1)*lda];
} else {
for (magma_int_t i=0; i < n-1; i++) {
d[i] = A[i*lda+nb]; // diag
e[i] = A[i*lda+nb-1]; // lower diag
}
d[n-1] = A[(n-1)*lda+nb];
}
#endif
return MAGMA_SUCCESS;
}
extern "C" magma_int_t
magma_zlobpcg(
magma_z_matrix A,
magma_z_solver_par *solver_par,
magma_z_preconditioner *precond_par,
magma_queue_t queue )
{
magma_int_t info = 0;
#define residualNorms(i,iter) ( residualNorms + (i) + (iter)*n )
#define SWAP(x, y) { pointer = x; x = y; y = pointer; }
#define hresidualNorms(i,iter) (hresidualNorms + (i) + (iter)*n )
#define gramA( m, n) (gramA + (m) + (n)*ldgram)
#define gramB( m, n) (gramB + (m) + (n)*ldgram)
#define gevectors(m, n) (gevectors + (m) + (n)*ldgram)
#define h_gramB( m, n) (h_gramB + (m) + (n)*ldgram)
#define magma_z_bspmv_tuned(m, n, alpha, A, X, beta, AX, queue) { \
magma_z_matrix x={Magma_CSR}, ax={Magma_CSR}; \
x.memory_location = Magma_DEV; x.num_rows = m; x.num_cols = n; x.major = MagmaColMajor; x.nnz = m*n; x.dval = X; x.storage_type = Magma_DENSE; \
ax.memory_location= Magma_DEV; ax.num_rows = m; ax.num_cols = n; ax.major = MagmaColMajor; ax.nnz = m*n; ax.dval = AX; ax.storage_type = Magma_DENSE; \
CHECK( magma_z_spmv(alpha, A, x, beta, ax, queue )); \
}
//**************************************************************
// Memory allocation for the eigenvectors, eigenvalues, and workspace
solver_par->solver = Magma_LOBPCG;
magma_int_t m = A.num_rows;
magma_int_t n = (solver_par->num_eigenvalues);
magmaDoubleComplex *blockX = solver_par->eigenvectors;
double *evalues = solver_par->eigenvalues;
solver_par->numiter = 0;
solver_par->spmv_count = 0;
magmaDoubleComplex *dwork=NULL, *hwork=NULL;
magmaDoubleComplex *blockP=NULL, *blockAP=NULL, *blockR=NULL, *blockAR=NULL, *blockAX=NULL, *blockW=NULL;
magmaDoubleComplex *gramA=NULL, *gramB=NULL, *gramM=NULL;
magmaDoubleComplex *gevectors=NULL, *h_gramB=NULL;
dwork = NULL;
hwork = NULL;
blockP = NULL;
blockR = NULL;
blockAP = NULL;
blockAR = NULL;
blockAX = NULL;
blockW = NULL;
gramA = NULL;
gramB = NULL;
gramM = NULL;
gevectors = NULL;
h_gramB = NULL;
magmaDoubleComplex *pointer, *origX = blockX;
double *eval_gpu=NULL;
magma_int_t iterationNumber, cBlockSize, restart = 1, iter;
//Chronometry
real_Double_t tempo1, tempo2, tempop1, tempop2;
magma_int_t lwork = max( 2*n+n*magma_get_dsytrd_nb(n),
1 + 6*3*n + 2* 3*n* 3*n);
magma_int_t *iwork={0}, liwork = 15*n+9;
magma_int_t gramDim, ldgram = 3*n, ikind = 3;
magmaDoubleComplex *hW={0};
// === Set solver parameters ===
double residualTolerance = solver_par->rtol;
magma_int_t maxIterations = solver_par->maxiter;
double tmp;
double r0=0; // set in 1st iteration
// === Set some constants & defaults ===
magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
double *residualNorms={0}, *condestGhistory={0}, condestG={0};
double *gevalues={0};
magma_int_t *activeMask={0};
double *hresidualNorms={0};
#ifdef COMPLEX
double *rwork={0};
magma_int_t lrwork = 1 + 5*(3*n) + 2*(3*n)*(3*n);
CHECK( magma_dmalloc_cpu(&rwork, lrwork));
#endif
CHECK( magma_zmalloc_pinned( &hwork , lwork ));
CHECK( magma_zmalloc( &blockAX , m*n ));
CHECK( magma_zmalloc( &blockAR , m*n ));
CHECK( magma_zmalloc( &blockAP , m*n ));
CHECK( magma_zmalloc( &blockR , m*n ));
CHECK( magma_zmalloc( &blockP , m*n ));
CHECK( magma_zmalloc( &blockW , m*n ));
CHECK( magma_zmalloc( &dwork , m*n ));
CHECK( magma_dmalloc( &eval_gpu , 3*n ));
//**********************************************************+
// === Check some parameters for possible quick exit ===
solver_par->info = MAGMA_SUCCESS;
if (m < 2)
info = MAGMA_DIVERGENCE;
else if (n > m)
info = MAGMA_SLOW_CONVERGENCE;
if (solver_par->info != 0) {
magma_xerbla( __func__, -(info) );
goto cleanup;
}
solver_par->info = info; // local info variable;
// === Allocate GPU memory for the residual norms' history ===
CHECK( magma_dmalloc(&residualNorms, (maxIterations+1) * n));
CHECK( magma_malloc( (void **)&activeMask, (n+1) * sizeof(magma_int_t) ));
// === Allocate CPU work space ===
CHECK( magma_dmalloc_cpu(&condestGhistory, maxIterations+1));
CHECK( magma_dmalloc_cpu(&gevalues, 3 * n));
CHECK( magma_malloc_cpu((void **)&iwork, liwork * sizeof(magma_int_t)));
CHECK( magma_zmalloc_pinned(&hW, n*n));
CHECK( magma_zmalloc_pinned(&gevectors, 9*n*n));
CHECK( magma_zmalloc_pinned(&h_gramB , 9*n*n));
// === Allocate GPU workspace ===
CHECK( magma_zmalloc(&gramM, n * n));
CHECK( magma_zmalloc(&gramA, 9 * n * n));
CHECK( magma_zmalloc(&gramB, 9 * n * n));
// === Set activemask to one ===
for(magma_int_t k =0; k<n; k++){
iwork[k]=1;
}
magma_setmatrix(n, 1, sizeof(magma_int_t), iwork, n , activeMask, n, queue);
#if defined(PRECISION_s)
ikind = 3;
#endif
// === Make the initial vectors orthonormal ===
magma_zgegqr_gpu(ikind, m, n, blockX, m, dwork, hwork, &info );
//magma_zorthomgs( m, n, blockX, queue );
magma_z_bspmv_tuned(m, n, c_one, A, blockX, c_zero, blockAX, queue );
solver_par->spmv_count++;
// === Compute the Gram matrix = (X, AX) & its eigenstates ===
magma_zgemm( MagmaConjTrans, MagmaNoTrans, n, n, m,
c_one, blockX, m, blockAX, m, c_zero, gramM, n, queue );
magma_zheevd_gpu( MagmaVec, MagmaUpper,
n, gramM, n, evalues, hW, n, hwork, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork, &info );
// === Update X = X * evectors ===
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockX, m, gramM, n, c_zero, blockW, m, queue );
SWAP(blockW, blockX);
// === Update AX = AX * evectors ===
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockAX, m, gramM, n, c_zero, blockW, m, queue );
SWAP(blockW, blockAX);
condestGhistory[1] = 7.82;
tempo1 = magma_sync_wtime( queue );
// === Main LOBPCG loop ============================================================
for(iterationNumber = 1; iterationNumber < maxIterations; iterationNumber++)
{
// === compute the residuals (R = Ax - x evalues )
magmablas_zlacpy( MagmaFull, m, n, blockAX, m, blockR, m, queue );
/*
for(magma_int_t i=0; i<n; i++) {
magma_zaxpy( m, MAGMA_Z_MAKE(-evalues[i],0), blockX+i*m, 1, blockR+i*m, 1, queue );
}
*/
magma_dsetmatrix( 3*n, 1, evalues, 3*n, eval_gpu, 3*n, queue );
CHECK( magma_zlobpcg_res( m, n, eval_gpu, blockX, blockR, eval_gpu, queue ));
magmablas_dznrm2_cols( m, n, blockR, m, residualNorms(0, iterationNumber), queue );
// === remove the residuals corresponding to already converged evectors
CHECK( magma_zcompact(m, n, blockR, m,
residualNorms(0, iterationNumber), residualTolerance,
activeMask, &cBlockSize, queue ));
if (cBlockSize == 0)
break;
// === apply a preconditioner P to the active residulas: R_new = P R_old
// === for now set P to be identity (no preconditioner => nothing to be done )
//magmablas_zlacpy( MagmaFull, m, cBlockSize, blockR, m, blockW, m, queue );
//SWAP(blockW, blockR);
// preconditioner
magma_z_matrix bWv={Magma_CSR}, bRv={Magma_CSR};
bWv.memory_location = Magma_DEV; bWv.num_rows = m; bWv.num_cols = cBlockSize; bWv.major = MagmaColMajor; bWv.nnz = m*cBlockSize; bWv.dval = blockW;
bRv.memory_location = Magma_DEV; bRv.num_rows = m; bRv.num_cols = cBlockSize; bRv.major = MagmaColMajor; bRv.nnz = m*cBlockSize; bRv.dval = blockR;
tempop1 = magma_sync_wtime( queue );
CHECK( magma_z_applyprecond_left( MagmaNoTrans, A, bRv, &bWv, precond_par, queue ));
CHECK( magma_z_applyprecond_right( MagmaNoTrans, A, bWv, &bRv, precond_par, queue ));
tempop2 = magma_sync_wtime( queue );
precond_par->runtime += tempop2-tempop1;
// === make the preconditioned residuals orthogonal to X
if( precond_par->solver != Magma_NONE){
magma_zgemm( MagmaConjTrans, MagmaNoTrans, n, cBlockSize, m,
c_one, blockX, m, blockR, m, c_zero, gramB(0,0), ldgram, queue );
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, cBlockSize, n,
c_neg_one, blockX, m, gramB(0,0), ldgram, c_one, blockR, m, queue );
}
// === make the active preconditioned residuals orthonormal
magma_zgegqr_gpu(ikind, m, cBlockSize, blockR, m, dwork, hwork, &info );
#if defined(PRECISION_s)
// re-orthogonalization
SWAP(blockX, dwork);
magma_zgegqr_gpu(ikind, m, cBlockSize, blockR, m, dwork, hwork, &info );
#endif
//magma_zorthomgs( m, cBlockSize, blockR, queue );
// === compute AR
magma_z_bspmv_tuned(m, cBlockSize, c_one, A, blockR, c_zero, blockAR, queue );
solver_par->spmv_count++;
if (!restart) {
// === compact P & AP as well
CHECK( magma_zcompactActive(m, n, blockP, m, activeMask, queue ));
CHECK( magma_zcompactActive(m, n, blockAP, m, activeMask, queue ));
/*
// === make P orthogonal to X ?
magma_zgemm( MagmaConjTrans, MagmaNoTrans, n, cBlockSize, m,
c_one, blockX, m, blockP, m, c_zero, gramB(0,0), ldgram, queue );
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, cBlockSize, n,
c_neg_one, blockX, m, gramB(0,0), ldgram, c_one, blockP, m, queue );
// === make P orthogonal to R ?
magma_zgemm( MagmaConjTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockR, m, blockP, m, c_zero, gramB(0,0), ldgram, queue );
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, cBlockSize, cBlockSize,
c_neg_one, blockR, m, gramB(0,0), ldgram, c_one, blockP, m, queue );
*/
// === Make P orthonormal & properly change AP (without multiplication by A)
magma_zgegqr_gpu(ikind, m, cBlockSize, blockP, m, dwork, hwork, &info );
#if defined(PRECISION_s)
// re-orthogonalization
SWAP(blockX, dwork);
magma_zgegqr_gpu(ikind, m, cBlockSize, blockP, m, dwork, hwork, &info );
#endif
//magma_zorthomgs( m, cBlockSize, blockP, queue );
//magma_z_bspmv_tuned(m, cBlockSize, c_one, A, blockP, c_zero, blockAP, queue );
magma_zsetmatrix( cBlockSize, cBlockSize, hwork, cBlockSize, dwork, cBlockSize, queue );
// replacement according to Stan
#if defined(PRECISION_s) || defined(PRECISION_d)
magmablas_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
m, cBlockSize, c_one, dwork, cBlockSize, blockAP, m, queue );
#else
magma_ztrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
m, cBlockSize, c_one, dwork, cBlockSize, blockAP, m, queue );
#endif
}
iter = max( 1, iterationNumber - 10 - int(log(1.*cBlockSize)) );
double condestGmean = 0.;
for(magma_int_t i = 0; i<iterationNumber-iter+1; i++){
condestGmean += condestGhistory[i];
}
condestGmean = condestGmean / (iterationNumber-iter+1);
if (restart)
gramDim = n+cBlockSize;
else
gramDim = n+2*cBlockSize;
/* --- The Raileight-Ritz method for [X R P] -----------------------
[ X R P ]' [AX AR AP] y = evalues [ X R P ]' [ X R P ], i.e.,
GramA GramB
/ X'AX X'AR X'AP \ / X'X X'R X'P \
| R'AX R'AR R'AP | y = evalues | R'X R'R R'P |
\ P'AX P'AR P'AP / \ P'X P'R P'P /
----------------------------------------------------------------- */
// === assemble GramB; first, set it to I
magmablas_zlaset( MagmaFull, ldgram, ldgram, c_zero, c_one, gramB, ldgram, queue ); // identity
if (!restart) {
magma_zgemm( MagmaConjTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockP, m, blockX, m, c_zero, gramB(n+cBlockSize,0), ldgram, queue );
magma_zgemm( MagmaConjTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockP, m, blockR, m, c_zero, gramB(n+cBlockSize,n), ldgram, queue );
}
magma_zgemm( MagmaConjTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockR, m, blockX, m, c_zero, gramB(n,0), ldgram, queue );
// === get GramB from the GPU to the CPU and compute its eigenvalues only
magma_zgetmatrix( gramDim, gramDim, gramB, ldgram, h_gramB, ldgram, queue );
lapackf77_zheev("N", "L", &gramDim, h_gramB, &ldgram, gevalues,
hwork, &lwork,
#ifdef COMPLEX
rwork,
#endif
&info);
// === check stability criteria if we need to restart
condestG = log10( gevalues[gramDim-1]/gevalues[0] ) + 1.;
if ((condestG/condestGmean>2 && condestG>2) || condestG>8) {
// Steepest descent restart for stability
restart=1;
printf("restart at step #%d\n", int(iterationNumber));
}
// === assemble GramA; first, set it to I
magmablas_zlaset( MagmaFull, ldgram, ldgram, c_zero, c_one, gramA, ldgram, queue ); // identity
magma_zgemm( MagmaConjTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockR, m, blockAX, m, c_zero, gramA(n,0), ldgram, queue );
magma_zgemm( MagmaConjTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockR, m, blockAR, m, c_zero, gramA(n,n), ldgram, queue );
if (!restart) {
magma_zgemm( MagmaConjTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockP, m, blockAX, m, c_zero,
gramA(n+cBlockSize,0), ldgram, queue );
magma_zgemm( MagmaConjTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockP, m, blockAR, m, c_zero,
gramA(n+cBlockSize,n), ldgram, queue );
magma_zgemm( MagmaConjTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockP, m, blockAP, m, c_zero,
gramA(n+cBlockSize,n+cBlockSize), ldgram, queue );
}
/*
// === Compute X' AX or just use the eigenvalues below ?
magma_zgemm( MagmaConjTrans, MagmaNoTrans, n, n, m,
c_one, blockX, m, blockAX, m, c_zero,
gramA(0,0), ldgram, queue );
*/
if (restart==0) {
magma_zgetmatrix( gramDim, gramDim, gramA, ldgram, gevectors, ldgram, queue );
}
else {
gramDim = n+cBlockSize;
magma_zgetmatrix( gramDim, gramDim, gramA, ldgram, gevectors, ldgram, queue );
}
for(magma_int_t k=0; k<n; k++)
*gevectors(k,k) = MAGMA_Z_MAKE(evalues[k], 0);
// === the previous eigensolver destroyed what is in h_gramB => must copy it again
magma_zgetmatrix( gramDim, gramDim, gramB, ldgram, h_gramB, ldgram, queue );
magma_int_t itype = 1;
lapackf77_zhegvd(&itype, "V", "L", &gramDim,
gevectors, &ldgram, h_gramB, &ldgram,
gevalues, hwork, &lwork,
#ifdef COMPLEX
rwork, &lrwork,
#endif
iwork, &liwork, &info);
for(magma_int_t k =0; k<n; k++)
evalues[k] = gevalues[k];
// === copy back the result to gramA on the GPU and use it for the updates
magma_zsetmatrix( gramDim, gramDim, gevectors, ldgram, gramA, ldgram, queue );
if (restart == 0) {
// === contribution from P to the new X (in new search direction P)
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockP, m, gramA(n+cBlockSize,0), ldgram, c_zero, dwork, m, queue );
SWAP(dwork, blockP);
// === contribution from R to the new X (in new search direction P)
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockR, m, gramA(n,0), ldgram, c_one, blockP, m, queue );
// === corresponding contribution from AP to the new AX (in AP)
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockAP, m, gramA(n+cBlockSize,0), ldgram, c_zero, dwork, m, queue );
SWAP(dwork, blockAP);
// === corresponding contribution from AR to the new AX (in AP)
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockAR, m, gramA(n,0), ldgram, c_one, blockAP, m, queue );
}
else {
// === contribution from R (only) to the new X
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockR, m, gramA(n,0), ldgram, c_zero, blockP, m, queue );
// === corresponding contribution from AR (only) to the new AX
magma_zgemm( MagmaNoTrans, MagmaNoTrans,m, n, cBlockSize,
c_one, blockAR, m, gramA(n,0), ldgram, c_zero, blockAP, m, queue );
}
// === contribution from old X to the new X + the new search direction P
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockX, m, gramA, ldgram, c_zero, dwork, m, queue );
SWAP(dwork, blockX);
//magma_zaxpy( m*n, c_one, blockP, 1, blockX, 1, queue );
CHECK( magma_zlobpcg_maxpy( m, n, blockP, blockX, queue ));
// === corresponding contribution from old AX to new AX + AP
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockAX, m, gramA, ldgram, c_zero, dwork, m, queue );
SWAP(dwork, blockAX);
//magma_zaxpy( m*n, c_one, blockAP, 1, blockAX, 1, queue );
CHECK( magma_zlobpcg_maxpy( m, n, blockAP, blockAX, queue ));
condestGhistory[iterationNumber+1]=condestG;
magma_dgetmatrix( 1, 1, residualNorms(0, iterationNumber), 1, &tmp, 1, queue );
if ( iterationNumber == 1 ) {
solver_par->init_res = tmp;
r0 = tmp * solver_par->rtol;
if ( r0 < ATOLERANCE )
r0 = ATOLERANCE;
}
solver_par->final_res = tmp;
if ( tmp < r0 ) {
break;
}
if (cBlockSize == 0) {
break;
}
if ( solver_par->verbose!=0 ) {
if ( iterationNumber%solver_par->verbose == 0 ) {
// double res;
// magma_zgetmatrix( 1, 1,
// (magmaDoubleComplex*)residualNorms(0, iterationNumber), 1,
// (magmaDoubleComplex*)&res, 1, queue );
//
// printf("Iteration %4d, CBS %4d, Residual: %10.7f\n",
// iterationNumber, cBlockSize, res);
printf("%4d-%2d ", int(iterationNumber), int(cBlockSize));
magma_dprint_gpu(1, n, residualNorms(0, iterationNumber), 1);
}
}
restart = 0;
} // === end for iterationNumber = 1,maxIterations =======================
// fill solver info
tempo2 = magma_sync_wtime( queue );
solver_par->runtime = (real_Double_t) tempo2-tempo1;
solver_par->numiter = iterationNumber;
if ( solver_par->numiter < solver_par->maxiter) {
info = MAGMA_SUCCESS;
} else if ( solver_par->init_res > solver_par->final_res )
info = MAGMA_SLOW_CONVERGENCE;
else
info = MAGMA_DIVERGENCE;
// =============================================================================
// === postprocessing;
// =============================================================================
// === compute the real AX and corresponding eigenvalues
magma_z_bspmv_tuned(m, n, c_one, A, blockX, c_zero, blockAX, queue );
magma_zgemm( MagmaConjTrans, MagmaNoTrans, n, n, m,
c_one, blockX, m, blockAX, m, c_zero, gramM, n, queue );
magma_zheevd_gpu( MagmaVec, MagmaUpper,
n, gramM, n, gevalues, dwork, n, hwork, lwork,
#ifdef COMPLEX
rwork, lrwork,
#endif
iwork, liwork, &info );
for(magma_int_t k =0; k<n; k++)
evalues[k] = gevalues[k];
// === update X = X * evectors
SWAP(blockX, dwork);
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, dwork, m, gramM, n, c_zero, blockX, m, queue );
// === update AX = AX * evectors to compute the final residual
SWAP(blockAX, dwork);
magma_zgemm( MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, dwork, m, gramM, n, c_zero, blockAX, m, queue );
// === compute R = AX - evalues X
magmablas_zlacpy( MagmaFull, m, n, blockAX, m, blockR, m, queue );
for(magma_int_t i=0; i<n; i++)
magma_zaxpy( m, MAGMA_Z_MAKE(-evalues[i], 0), blockX+i*m, 1, blockR+i*m, 1, queue );
// === residualNorms[iterationNumber] = || R ||
magmablas_dznrm2_cols( m, n, blockR, m, residualNorms(0, iterationNumber), queue );
// === restore blockX if needed
if (blockX != origX)
magmablas_zlacpy( MagmaFull, m, n, blockX, m, origX, m, queue );
printf("Eigenvalues:\n");
for(magma_int_t i =0; i<n; i++)
printf("%e ", evalues[i]);
printf("\n\n");
printf("Final residuals:\n");
magma_dprint_gpu(1, n, residualNorms(0, iterationNumber), 1);
printf("\n\n");
//=== Prmagma_int_t residual history in a file for plotting ====
CHECK( magma_dmalloc_cpu(&hresidualNorms, (iterationNumber+1) * n));
magma_dgetmatrix( n, iterationNumber,
residualNorms, n,
hresidualNorms, n, queue );
solver_par->iter_res = *hresidualNorms(0, iterationNumber-1);
printf("Residuals are stored in file residualNorms\n");
printf("Plot the residuals using: myplot \n");
FILE *residuals_file;
residuals_file = fopen("residualNorms", "w");
for(magma_int_t i =1; i<iterationNumber; i++) {
for(magma_int_t j = 0; j<n; j++)
fprintf(residuals_file, "%f ", *hresidualNorms(j,i));
fprintf(residuals_file, "\n");
}
fclose(residuals_file);
cleanup:
magma_free_cpu(hresidualNorms);
// === free work space
magma_free( residualNorms );
magma_free_cpu( condestGhistory );
magma_free_cpu( gevalues );
magma_free_cpu( iwork );
magma_free_pinned( hW );
magma_free_pinned( gevectors );
magma_free_pinned( h_gramB );
magma_free( gramM );
magma_free( gramA );
magma_free( gramB );
magma_free( activeMask );
if (blockX != (solver_par->eigenvectors))
magma_free( blockX );
if (blockAX != (solver_par->eigenvectors))
magma_free( blockAX );
if (blockAR != (solver_par->eigenvectors))
magma_free( blockAR );
if (blockAP != (solver_par->eigenvectors))
magma_free( blockAP );
if (blockR != (solver_par->eigenvectors))
magma_free( blockR );
if (blockP != (solver_par->eigenvectors))
magma_free( blockP );
if (blockW != (solver_par->eigenvectors))
magma_free( blockW );
if (dwork != (solver_par->eigenvectors))
magma_free( dwork );
magma_free( eval_gpu );
magma_free_pinned( hwork );
#ifdef COMPLEX
magma_free_cpu( rwork );
rwork = NULL;
#endif
return info;
}
/**
Purpose
-------
ZPOTRF computes the Cholesky factorization of a complex Hermitian
positive definite matrix dA.
The factorization has the form
dA = U**H * U, if UPLO = MagmaUpper, or
dA = L * L**H, if UPLO = MagmaLower,
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
If the current stream is NULL, this version replaces it with a new
stream to overlap computation with communication.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of dA is stored;
- = MagmaLower: Lower triangle of dA is stored.
@param[in]
n INTEGER
The order of the matrix dA. N >= 0.
@param[in,out]
dA COMPLEX_16 array on the GPU, dimension (LDDA,N)
On entry, the Hermitian matrix dA. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of dA contains the upper
triangular part of the matrix dA, and the strictly lower
triangular part of dA is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of dA contains the lower
triangular part of the matrix dA, and the strictly upper
triangular part of dA is not referenced.
\n
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization dA = U**H * U or dA = L * L**H.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
@ingroup magma_zposv_comp
********************************************************************/
extern "C" magma_int_t
magma_zpotrf_batched(
magma_uplo_t uplo, magma_int_t n,
magmaDoubleComplex **dA_array, magma_int_t ldda,
magma_int_t *info_array, magma_int_t batchCount, magma_queue_t queue)
{
#define A(i_, j_) (A + (i_) + (j_)*ldda)
double d_alpha = -1.0;
double d_beta = 1.0;
cudaMemset(info_array, 0, batchCount*sizeof(magma_int_t));
magma_int_t arginfo = 0;
if ( uplo != MagmaUpper && uplo != MagmaLower) {
arginfo = -1;
} else if (n < 0) {
arginfo = -2;
} else if (ldda < max(1,n)) {
arginfo = -4;
}
if (arginfo != 0) {
magma_xerbla( __func__, -(arginfo) );
return arginfo;
}
// Quick return if possible
if (n == 0) {
return arginfo;
}
if( n > 2048 ){
printf("=========================================================================================\n");
printf(" WARNING batched routines are designed for small sizes it might be better to use the\n Native/Hybrid classical routines if you want performance\n");
printf("=========================================================================================\n");
}
magma_int_t j, k, ib;
magma_int_t nb = POTRF_NB;
magma_int_t gemm_crossover = 127;//nb > 32 ? 127 : 160;
#if defined(USE_CUOPT)
cublasHandle_t myhandle;
cublasCreate_v2(&myhandle);
#else
cublasHandle_t myhandle=NULL;
#endif
magmaDoubleComplex **dA_displ = NULL;
magmaDoubleComplex **dW0_displ = NULL;
magmaDoubleComplex **dW1_displ = NULL;
magmaDoubleComplex **dW2_displ = NULL;
magmaDoubleComplex **dW3_displ = NULL;
magmaDoubleComplex **dW4_displ = NULL;
magmaDoubleComplex **dinvA_array = NULL;
magmaDoubleComplex **dwork_array = NULL;
magma_malloc((void**)&dA_displ, batchCount * sizeof(*dA_displ));
magma_malloc((void**)&dW0_displ, batchCount * sizeof(*dW0_displ));
magma_malloc((void**)&dW1_displ, batchCount * sizeof(*dW1_displ));
magma_malloc((void**)&dW2_displ, batchCount * sizeof(*dW2_displ));
magma_malloc((void**)&dW3_displ, batchCount * sizeof(*dW3_displ));
magma_malloc((void**)&dW4_displ, batchCount * sizeof(*dW4_displ));
magma_malloc((void**)&dinvA_array, batchCount * sizeof(*dinvA_array));
magma_malloc((void**)&dwork_array, batchCount * sizeof(*dwork_array));
magma_int_t invA_msize = ((n+TRI_NB-1)/TRI_NB)*TRI_NB*TRI_NB;
magma_int_t dwork_msize = n*nb;
magmaDoubleComplex* dinvA = NULL;
magmaDoubleComplex* dwork = NULL;// dinvA and dwork are workspace in ztrsm
magmaDoubleComplex **cpuAarray = NULL;
magma_zmalloc( &dinvA, invA_msize * batchCount);
magma_zmalloc( &dwork, dwork_msize * batchCount );
magma_malloc_cpu((void**) &cpuAarray, batchCount*sizeof(magmaDoubleComplex*));
/* check allocation */
if ( dA_displ == NULL || dW0_displ == NULL || dW1_displ == NULL || dW2_displ == NULL ||
dW3_displ == NULL || dW4_displ == NULL || dinvA_array == NULL || dwork_array == NULL ||
dinvA == NULL || dwork == NULL || cpuAarray == NULL ) {
magma_free(dA_displ);
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dinvA_array);
magma_free(dwork_array);
magma_free( dinvA );
magma_free( dwork );
free(cpuAarray);
magma_int_t info = MAGMA_ERR_DEVICE_ALLOC;
magma_xerbla( __func__, -(info) );
return info;
}
magmablas_zlaset_q(MagmaFull, invA_msize, batchCount, MAGMA_Z_ZERO, MAGMA_Z_ZERO, dinvA, invA_msize, queue);
magmablas_zlaset_q(MagmaFull, dwork_msize, batchCount, MAGMA_Z_ZERO, MAGMA_Z_ZERO, dwork, dwork_msize, queue);
zset_pointer(dwork_array, dwork, 1, 0, 0, dwork_msize, batchCount, queue);
zset_pointer(dinvA_array, dinvA, TRI_NB, 0, 0, invA_msize, batchCount, queue);
magma_queue_t cstream;
magmablasGetKernelStream(&cstream);
magma_int_t streamid;
const magma_int_t nbstreams=32;
magma_queue_t stream[nbstreams];
for(k=0; k<nbstreams; k++){
magma_queue_create( &stream[k] );
}
magma_getvector( batchCount, sizeof(magmaDoubleComplex*), dA_array, 1, cpuAarray, 1);
magmablasSetKernelStream(NULL);
if (uplo == MagmaUpper) {
printf("Upper side is unavailable \n");
goto fin;
}
else {
for(j = 0; j < n; j+=nb) {
ib = min(nb, n-j);
#if 1
//===============================================
// panel factorization
//===============================================
magma_zdisplace_pointers(dA_displ, dA_array, ldda, j, j, batchCount, queue);
zset_pointer(dwork_array, dwork, 1, 0, 0, dwork_msize, batchCount, queue);
zset_pointer(dinvA_array, dinvA, TRI_NB, 0, 0, invA_msize, batchCount, queue);
#if 0
arginfo = magma_zpotrf_panel_batched(
uplo, n-j, ib,
dA_displ, ldda,
dwork_array, dwork_msize,
dinvA_array, invA_msize,
dW0_displ, dW1_displ, dW2_displ,
dW3_displ, dW4_displ,
info_array, j, batchCount, myhandle);
#else
//arginfo = magma_zpotrf_rectile_batched(
arginfo = magma_zpotrf_recpanel_batched(
uplo, n-j, ib, 32,
dA_displ, ldda,
dwork_array, dwork_msize,
dinvA_array, invA_msize,
dW0_displ, dW1_displ, dW2_displ,
dW3_displ, dW4_displ,
info_array, j, batchCount, myhandle, queue);
#endif
if(arginfo != 0 ) goto fin;
//===============================================
// end of panel
//===============================================
#endif
#if 1
//real_Double_t gpu_time;
//gpu_time = magma_sync_wtime(NULL);
if( (n-j-ib) > 0){
if( (n-j-ib) > gemm_crossover)
{
//-------------------------------------------
// USE STREAM HERK
//-------------------------------------------
// since it use different stream I need to wait the panel.
// But since the code use the NULL stream everywhere,
// so I don't need it, because the NULL stream do the sync by itself
//magma_queue_sync(NULL);
/* you must know the matrix layout inorder to do it */
for(k=0; k<batchCount; k++)
{
streamid = k%nbstreams;
magmablasSetKernelStream(stream[streamid]);
// call herk, class zherk must call cpu pointer
magma_zherk(MagmaLower, MagmaNoTrans, n-j-ib, ib,
d_alpha,
(const magmaDoubleComplex*) cpuAarray[k] + j+ib+j*ldda, ldda,
d_beta,
cpuAarray[k] + j+ib+(j+ib)*ldda, ldda);
}
// need to synchronise to be sure that panel do not start before
// finishing the update at least of the next panel
// BUT no need for it as soon as the other portion of the code
// use the NULL stream which do the sync by itself
//magma_device_sync();
magmablasSetKernelStream(NULL);
}
else
{
//-------------------------------------------
// USE BATCHED GEMM(which is a HERK in fact, since it only access the lower part)
//-------------------------------------------
magma_zdisplace_pointers(dA_displ, dA_array, ldda, j+ib, j, batchCount, queue);
magma_zdisplace_pointers(dW1_displ, dA_array, ldda, j+ib, j+ib, batchCount, queue);
magmablas_zherk_batched(uplo, MagmaNoTrans, n-j-ib, ib,
d_alpha, dA_displ, ldda,
d_beta, dW1_displ, ldda,
batchCount, queue);
}
}
//gpu_time = magma_sync_wtime(NULL) - gpu_time;
//real_Double_t flops = (n-j-ib) * (n-j-ib) * ib / 1e9 * batchCount;
//real_Double_t gpu_perf = flops / gpu_time;
//printf("Rows= %d, Colum=%d, herk time = %7.2fms, Gflops= %7.2f\n", n-j-ib, ib, gpu_time*1000, gpu_perf);
#endif
}
}
fin:
magma_queue_sync(NULL);
for(k=0; k<nbstreams; k++){
magma_queue_destroy( stream[k] );
}
magmablasSetKernelStream(cstream);
#if defined(USE_CUOPT)
cublasDestroy_v2(myhandle);
#endif
magma_free(dA_displ);
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dinvA_array);
magma_free(dwork_array);
magma_free( dinvA );
magma_free( dwork );
free(cpuAarray);
return arginfo;
}
/**
Purpose
-------
SGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
This is a batched version that factors batchCount M-by-N matrices in parallel.
dA, ipiv, and info become arrays with one entry per matrix.
Arguments
---------
@param[in]
m INTEGER
The number of rows of each matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of each matrix A. N >= 0.
@param[in,out]
dA_array Array of pointers, dimension (batchCount).
Each is a REAL array on the GPU, dimension (LDDA,N).
On entry, each pointer is an M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
@param[in]
ldda INTEGER
The leading dimension of each array A. LDDA >= max(1,M).
@param[out]
ipiv_array Array of pointers, dimension (batchCount), for corresponding matrices.
Each is an INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
@param[out]
info_array Array of INTEGERs, dimension (batchCount), for corresponding matrices.
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
- > 0: if INFO = i, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
@param[in]
batchCount INTEGER
The number of matrices to operate on.
@param[in]
queue magma_queue_t
Queue to execute in.
@ingroup magma_sgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_sgetrf_batched(
magma_int_t m, magma_int_t n,
float **dA_array,
magma_int_t ldda,
magma_int_t **ipiv_array,
magma_int_t *info_array,
magma_int_t batchCount, magma_queue_t queue)
{
#define A(i_, j_) (A + (i_) + (j_)*ldda)
magma_int_t min_mn = min(m, n);
cudaMemset(info_array, 0, batchCount*sizeof(magma_int_t));
/* Check arguments */
magma_int_t arginfo = 0;
if (m < 0)
arginfo = -1;
else if (n < 0)
arginfo = -2;
else if (ldda < max(1,m))
arginfo = -4;
if (arginfo != 0) {
magma_xerbla( __func__, -(arginfo) );
return arginfo;
}
/* Quick return if possible */
if (m == 0 || n == 0)
if (min_mn == 0 ) return arginfo;
if ( m > 2048 || n > 2048 ) {
#ifndef MAGMA_NOWARNING
printf("=========================================================================================\n");
printf(" WARNING batched routines are designed for small sizes it might be better to use the\n Native/Hybrid classical routines if you want performance\n");
printf("=========================================================================================\n");
#endif
}
//#define ENABLE_TIMER3
#if defined(ENABLE_TIMER3)
real_Double_t tall=0.0, tloop=0., talloc=0., tdalloc=0.;
tall = magma_sync_wtime(queue);
talloc = magma_sync_wtime(queue);
#endif
float c_neg_one = MAGMA_S_NEG_ONE;
float c_one = MAGMA_S_ONE;
magma_int_t nb, recnb, ib, i, k, pm, use_stream;
magma_get_sgetrf_batched_nbparam(n, &nb, &recnb);
magma_int_t **dipiv_displ = NULL;
float **dA_displ = NULL;
float **dW0_displ = NULL;
float **dW1_displ = NULL;
float **dW2_displ = NULL;
float **dW3_displ = NULL;
float **dW4_displ = NULL;
float **dinvA_array = NULL;
float **dwork_array = NULL;
magma_malloc((void**)&dipiv_displ, batchCount * sizeof(*dipiv_displ));
magma_malloc((void**)&dA_displ, batchCount * sizeof(*dA_displ));
magma_malloc((void**)&dW0_displ, batchCount * sizeof(*dW0_displ));
magma_malloc((void**)&dW1_displ, batchCount * sizeof(*dW1_displ));
magma_malloc((void**)&dW2_displ, batchCount * sizeof(*dW2_displ));
magma_malloc((void**)&dW3_displ, batchCount * sizeof(*dW3_displ));
magma_malloc((void**)&dW4_displ, batchCount * sizeof(*dW4_displ));
magma_malloc((void**)&dinvA_array, batchCount * sizeof(*dinvA_array));
magma_malloc((void**)&dwork_array, batchCount * sizeof(*dwork_array));
magma_int_t invA_msize = magma_roundup( n, TRI_NB )*TRI_NB;
magma_int_t dwork_msize = n*nb;
magma_int_t **pivinfo_array = NULL;
magma_int_t *pivinfo = NULL;
float* dinvA = NULL;
float* dwork = NULL; // dinvA and dwork are workspace in strsm
float **cpuAarray = NULL;
magma_smalloc( &dinvA, invA_msize * batchCount);
magma_smalloc( &dwork, dwork_msize * batchCount );
magma_malloc((void**)&pivinfo_array, batchCount * sizeof(*pivinfo_array));
magma_malloc((void**)&pivinfo, batchCount * m * sizeof(magma_int_t));
magma_malloc_cpu((void**) &cpuAarray, batchCount*sizeof(float*));
/* check allocation */
if ( dA_displ == NULL || dW0_displ == NULL || dW1_displ == NULL || dW2_displ == NULL ||
dW3_displ == NULL || dW4_displ == NULL || dinvA_array == NULL || dwork_array == NULL ||
dinvA == NULL || dwork == NULL || cpuAarray == NULL ||
dipiv_displ == NULL || pivinfo_array == NULL || pivinfo == NULL) {
magma_free(dA_displ);
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dinvA_array);
magma_free(dwork_array);
magma_free( dinvA );
magma_free( dwork );
magma_free_cpu(cpuAarray);
magma_free(dipiv_displ);
magma_free(pivinfo_array);
magma_free(pivinfo);
magma_int_t info = MAGMA_ERR_DEVICE_ALLOC;
magma_xerbla( __func__, -(info) );
return info;
}
magmablas_slaset_q( MagmaFull, invA_msize, batchCount, MAGMA_S_ZERO, MAGMA_S_ZERO, dinvA, invA_msize, queue );
magmablas_slaset_q( MagmaFull, dwork_msize, batchCount, MAGMA_S_ZERO, MAGMA_S_ZERO, dwork, dwork_msize, queue );
magma_sset_pointer( dwork_array, dwork, 1, 0, 0, dwork_msize, batchCount, queue );
magma_sset_pointer( dinvA_array, dinvA, TRI_NB, 0, 0, invA_msize, batchCount, queue );
magma_iset_pointer( pivinfo_array, pivinfo, 1, 0, 0, m, batchCount, queue );
magma_int_t streamid;
const magma_int_t nbstreams=10;
magma_queue_t queues[nbstreams];
for (i=0; i < nbstreams; i++) {
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[i] );
}
magma_getvector( batchCount, sizeof(float*), dA_array, 1, cpuAarray, 1, queue);
#if defined(ENABLE_TIMER3)
printf(" I am after malloc\n");
talloc = magma_sync_wtime(queue) - talloc;
tloop = magma_sync_wtime(queue);
#endif
for (i = 0; i < min_mn; i += nb)
{
ib = min(nb, min_mn-i);
pm = m-i;
magma_idisplace_pointers(dipiv_displ, ipiv_array, ldda, i, 0, batchCount, queue);
magma_sdisplace_pointers(dA_displ, dA_array, ldda, i, i, batchCount, queue);
//===============================================
// panel factorization
//===============================================
if (recnb == nb)
{
arginfo = magma_sgetf2_batched(
pm, ib,
dA_displ, ldda,
dW1_displ, dW2_displ, dW3_displ,
dipiv_displ,
info_array, i, batchCount, queue);
}
else {
arginfo = magma_sgetrf_recpanel_batched(
pm, ib, recnb,
dA_displ, ldda,
dipiv_displ, pivinfo_array,
dwork_array, nb,
dinvA_array, invA_msize,
dW0_displ, dW1_displ, dW2_displ,
dW3_displ, dW4_displ,
info_array, i,
batchCount, queue);
}
if (arginfo != 0 ) goto fin;
//===============================================
// end of panel
//===============================================
#define RUN_ALL
#ifdef RUN_ALL
// setup pivinfo before adjusting ipiv
setup_pivinfo_batched(pivinfo_array, dipiv_displ, pm, ib, batchCount, queue);
adjust_ipiv_batched(dipiv_displ, ib, i, batchCount, queue);
// stepinit_ipiv(pivinfo_array, pm, batchCount); // for debug and check swap, it create an ipiv
#if 0
slaswp_batched( i, dA_displ, ldda,
i, i+ib,
dipiv_displ, pivinfo_array, batchCount);
#else
magma_sdisplace_pointers(dA_displ, dA_array, ldda, i, 0, batchCount, queue);
magma_sdisplace_pointers(dW0_displ, dA_array, ldda, i, 0, batchCount, queue);
magma_slaswp_rowparallel_batched( i, dA_displ, ldda,
dW0_displ, ldda,
i, i+ib,
pivinfo_array, batchCount, queue );
#endif
if ( (i + ib) < n)
{
// swap right side and trsm
magma_sdisplace_pointers(dA_displ, dA_array, ldda, i, i+ib, batchCount, queue);
magma_sset_pointer( dwork_array, dwork, nb, 0, 0, dwork_msize, batchCount, queue ); // I don't think it is needed Azzam
magma_slaswp_rowparallel_batched( n-(i+ib), dA_displ, ldda,
dwork_array, nb,
i, i+ib,
pivinfo_array, batchCount, queue );
magma_sdisplace_pointers(dA_displ, dA_array, ldda, i, i, batchCount, queue);
magma_sdisplace_pointers(dW0_displ, dA_array, ldda, i, i+ib, batchCount, queue);
magmablas_strsm_outofplace_batched( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit, 1,
ib, n-i-ib,
MAGMA_S_ONE,
dA_displ, ldda, // dA
dwork_array, nb, // dB
dW0_displ, ldda, // dX
dinvA_array, invA_msize,
dW1_displ, dW2_displ,
dW3_displ, dW4_displ,
0, batchCount, queue );
if ( (i + ib) < m)
{
// if gemm size is > 160 use a streamed classical cublas gemm since it is faster
// the batched is faster only when M=N <= 160 for K40c
//-------------------------------------------
// USE STREAM GEMM
//-------------------------------------------
use_stream = magma_srecommend_cublas_gemm_stream(MagmaNoTrans, MagmaNoTrans, m-i-ib, n-i-ib, ib);
if (use_stream)
{
magma_queue_sync(queue);
for (k=0; k < batchCount; k++)
{
streamid = k%nbstreams;
magma_sgemm( MagmaNoTrans, MagmaNoTrans,
m-i-ib, n-i-ib, ib,
c_neg_one, cpuAarray[k] + (i+ib)+i*ldda, ldda,
cpuAarray[k] + i+(i+ib)*ldda, ldda,
c_one, cpuAarray[k] + (i+ib)+(i+ib)*ldda, ldda, queues[streamid] );
}
// need to synchronise to be sure that sgetf2 do not start before
// finishing the update at least of the next panel
// if queue is NULL, no need to sync
if ( queue != NULL ) {
for (magma_int_t s=0; s < nbstreams; s++)
magma_queue_sync(queues[s]);
}
}
//-------------------------------------------
// USE BATCHED GEMM
//-------------------------------------------
else
{
magma_sdisplace_pointers(dA_displ, dA_array, ldda, i+ib, i, batchCount, queue);
magma_sdisplace_pointers(dW1_displ, dA_array, ldda, i, i+ib, batchCount, queue);
magma_sdisplace_pointers(dW2_displ, dA_array, ldda, i+ib, i+ib, batchCount, queue);
//printf("caling batched dgemm %d %d %d \n", m-i-ib, n-i-ib, ib);
magma_sgemm_batched( MagmaNoTrans, MagmaNoTrans, m-i-ib, n-i-ib, ib,
c_neg_one, dA_displ, ldda,
dW1_displ, ldda,
c_one, dW2_displ, ldda,
batchCount, queue );
} // end of batched/streamed gemm
} // end of if ( (i + ib) < m)
} // end of if ( (i + ib) < n)
#endif
}// end of for
fin:
magma_queue_sync(queue);
#if defined(ENABLE_TIMER3)
tloop = magma_sync_wtime(queue) - tloop;
tdalloc = magma_sync_wtime(queue);
#endif
for (k=0; k < nbstreams; k++) {
magma_queue_destroy( queues[k] );
}
magma_free(dA_displ);
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dinvA_array);
magma_free(dwork_array);
magma_free( dinvA );
magma_free( dwork );
magma_free_cpu(cpuAarray);
magma_free(dipiv_displ);
magma_free(pivinfo_array);
magma_free(pivinfo);
#if defined(ENABLE_TIMER3)
tdalloc = magma_sync_wtime(queue) - tdalloc;
tall = magma_sync_wtime(queue) - tall;
printf("here is the timing from inside sgetrf_batched talloc: %10.5f tloop: %10.5f tdalloc: %10.5f tall: %10.5f sum: %10.5f\n", talloc, tloop, tdalloc, tall, talloc+tloop+tdalloc );
#endif
return arginfo;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing ztrsm
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cublas_perf, cublas_time, cpu_perf=0, cpu_time=0;
double magma_error, cublas_error, work[1];
magma_int_t M, N, info;
magma_int_t Ak;
magma_int_t sizeA, sizeB;
magma_int_t lda, ldb, ldda, lddb;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t *piv;
magma_err_t err;
magmaDoubleComplex *h_A, *h_B, *h_Bcublas, *h_Bmagma, *h_B1, *h_X1, *h_X2, *LU, *LUT;
magmaDoubleComplex *d_A, *d_B;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 0.29, -0.86 );
magma_opts opts;
parse_opts( argc, argv, &opts );
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"
"side = %c, uplo = %c, transA = %c, diag = %c \n", opts.side, opts.uplo, opts.transA, opts.diag );
printf(" M N MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) CPU Gflop/s (ms) MAGMA error CUBLAS error\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];
gflops = FLOPS_ZTRSM(opts.side, M, N) / 1e9;
if ( opts.side == MagmaLeft ) {
lda = M;
Ak = M;
} else {
lda = N;
Ak = N;
}
ldb = M;
ldda = ((lda+31)/32)*32;
lddb = ((ldb+31)/32)*32;
sizeA = lda*Ak;
sizeB = ldb*N;
TESTING_MALLOC( h_A, magmaDoubleComplex, lda*Ak );
TESTING_MALLOC( LU, magmaDoubleComplex, lda*Ak );
TESTING_MALLOC( LUT, magmaDoubleComplex, lda*Ak );
TESTING_MALLOC( h_B, magmaDoubleComplex, ldb*N );
TESTING_MALLOC( h_B1, magmaDoubleComplex, ldb*N );
TESTING_MALLOC( h_X1, magmaDoubleComplex, ldb*N );
TESTING_MALLOC( h_X2, magmaDoubleComplex, ldb*N );
TESTING_MALLOC( h_Bcublas, magmaDoubleComplex, ldb*N );
TESTING_MALLOC( h_Bmagma, magmaDoubleComplex, ldb*N );
TESTING_DEVALLOC( d_A, magmaDoubleComplex, ldda*Ak );
TESTING_DEVALLOC( d_B, magmaDoubleComplex, lddb*N );
/* Initialize the matrices */
lapackf77_zlarnv( &ione, ISEED, &sizeA, LU );
err = magma_malloc_cpu( (void**) &piv, Ak*sizeof(magma_int_t) ); assert( err == 0 );
lapackf77_zgetrf( &Ak, &Ak, LU, &lda, piv, &info );
int i, j;
for(i=0;i<Ak;i++){
for(j=0;j<Ak;j++){
LUT[j+i*lda] = LU[i+j*lda];
}
}
lapackf77_zlacpy(MagmaUpperStr, &Ak, &Ak, LUT, &lda, LU, &lda);
if(opts.uplo == MagmaLower){
lapackf77_zlacpy(MagmaLowerStr, &Ak, &Ak, LU, &lda, h_A, &lda);
}else{
lapackf77_zlacpy(MagmaUpperStr, &Ak, &Ak, LU, &lda, h_A, &lda);
}
lapackf77_zlarnv( &ione, ISEED, &sizeB, h_B );
memcpy(h_B1, h_B, sizeB*sizeof(magmaDoubleComplex));
/* =====================================================================
Performs operation using MAGMA-BLAS
=================================================================== */
magma_zsetmatrix( Ak, Ak, h_A, lda, d_A, ldda );
magma_zsetmatrix( M, N, h_B, ldb, d_B, lddb );
magma_time = magma_sync_wtime( NULL );
magmablas_ztrsm( opts.side, opts.uplo, opts.transA, opts.diag,
M, N,
alpha, d_A, ldda,
d_B, lddb );
magma_time = magma_sync_wtime( NULL ) - magma_time;
magma_perf = gflops / magma_time;
magma_zgetmatrix( M, N, d_B, lddb, h_Bmagma, ldb );
/* =====================================================================
Performs operation using CUDA-BLAS
=================================================================== */
magma_zsetmatrix( M, N, h_B, ldb, d_B, lddb );
cublas_time = magma_sync_wtime( NULL );
cublasZtrsm( opts.side, opts.uplo, opts.transA, opts.diag,
M, N,
alpha, d_A, ldda,
d_B, lddb );
cublas_time = magma_sync_wtime( NULL ) - cublas_time;
cublas_perf = gflops / cublas_time;
magma_zgetmatrix( M, N, d_B, lddb, h_Bcublas, ldb );
/* =====================================================================
Performs operation using CPU BLAS
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
blasf77_ztrsm( &opts.side, &opts.uplo, &opts.transA, &opts.diag,
&M, &N,
&alpha, h_A, &lda,
h_B, &ldb );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
}
/* =====================================================================
Check the result
=================================================================== */
// ||b - Ax|| / (||A||*||x||)
memcpy(h_X1, h_Bmagma, sizeB*sizeof(magmaDoubleComplex));
magmaDoubleComplex alpha2 = MAGMA_Z_DIV( c_one, alpha );
blasf77_ztrmm( &opts.side, &opts.uplo, &opts.transA, &opts.diag,
&M, &N,
&alpha2, h_A, &lda,
h_X1, &ldb );
blasf77_zaxpy( &sizeB, &c_neg_one, h_B1, &ione, h_X1, &ione );
double norm1 = lapackf77_zlange( "M", &M, &N, h_X1, &ldb, work );
double normx = lapackf77_zlange( "M", &M, &N, h_Bmagma, &ldb, work );
double normA = lapackf77_zlange( "M", &Ak, &Ak, h_A, &lda, work );
magma_error = norm1/(normx*normA);
memcpy(h_X2, h_Bcublas, sizeB*sizeof(magmaDoubleComplex));
blasf77_ztrmm( &opts.side, &opts.uplo, &opts.transA, &opts.diag,
&M, &N,
&alpha2, h_A, &lda,
h_X2, &ldb );
blasf77_zaxpy( &sizeB, &c_neg_one, h_B1, &ione, h_X2, &ione );
norm1 = lapackf77_zlange( "M", &M, &N, h_X2, &ldb, work );
normx = lapackf77_zlange( "M", &M, &N, h_Bcublas, &ldb, work );
normA = lapackf77_zlange( "M", &Ak, &Ak, h_A, &lda, work );
cublas_error = norm1/(normx*normA);
if ( opts.lapack ) {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, cublas_error );
}
else {
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) --- ( --- ) %8.2e %8.2e\n",
(int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
magma_error, cublas_error );
}
TESTING_FREE( h_A );
TESTING_FREE( LU );
TESTING_FREE( LUT );
TESTING_FREE( h_B );
TESTING_FREE( h_Bcublas );
TESTING_FREE( h_Bmagma );
TESTING_FREE( h_B1 );
TESTING_FREE( h_X1 );
TESTING_FREE( h_X2 );
TESTING_DEVFREE( d_A );
TESTING_DEVFREE( d_B );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return 0;
}
/**
Purpose
-------
DGETRF computes an LU factorization of a general M-by-N matrix A
using partial pivoting with row interchanges.
The factorization has the form
A = P * L * U
where P is a permutation matrix, L is lower triangular with unit
diagonal elements (lower trapezoidal if m > n), and U is upper
triangular (upper trapezoidal if m < n).
This is the right-looking Level 3 BLAS version of the algorithm.
If the current stream is NULL, this version replaces it with a new
stream to overlap computation with communication.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix A. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix A. N >= 0.
@param[in,out]
dA DOUBLE_PRECISION array on the GPU, dimension (LDDA,N).
On entry, the M-by-N matrix to be factored.
On exit, the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDDA >= max(1,M).
@param[out]
ipiv INTEGER array, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was interchanged with row IPIV(i).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
or another error occured, such as memory allocation failed.
- > 0: if INFO = i, U(i,i) is exactly zero. The factorization
has been completed, but the factor U is exactly
singular, and division by zero will occur if it is used
to solve a system of equations.
@ingroup magma_dgesv_comp
********************************************************************/
extern "C" magma_int_t
magma_dgetrf_batched(
magma_int_t m, magma_int_t n,
double **dA_array,
magma_int_t ldda,
magma_int_t **ipiv_array,
magma_int_t *info_array,
magma_int_t batchCount, magma_queue_t queue)
{
#define A(i_, j_) (A + (i_) + (j_)*ldda)
magma_int_t min_mn = min(m, n);
cudaMemset(info_array, 0, batchCount*sizeof(magma_int_t));
/* Check arguments */
magma_int_t arginfo = 0;
if (m < 0)
arginfo = -1;
else if (n < 0)
arginfo = -2;
else if (ldda < max(1,m))
arginfo = -4;
if (arginfo != 0) {
magma_xerbla( __func__, -(arginfo) );
return arginfo;
}
/* Quick return if possible */
if (m == 0 || n == 0)
if(min_mn == 0 ) return arginfo;
if( m > 2048 || n > 2048 ){
printf("=========================================================================================\n");
printf(" WARNING batched routines are designed for small sizes it might be better to use the\n Native/Hybrid classical routines if you want performance\n");
printf("=========================================================================================\n");
}
//#define ENABLE_TIMER3
#if defined(ENABLE_TIMER3)
real_Double_t tall=0.0, tloop=0., talloc=0., tdalloc=0.;
tall = magma_sync_wtime(0);
talloc = magma_sync_wtime(0);
#endif
double neg_one = MAGMA_D_NEG_ONE;
double one = MAGMA_D_ONE;
magma_int_t ib, i, k, pm;
magma_int_t nb = BATRF_NB;
magma_int_t gemm_crossover = nb > 32 ? 127 : 160;
// magma_int_t gemm_crossover = n;// use only stream gemm
#if defined(USE_CUOPT)
cublasHandle_t myhandle;
cublasCreate_v2(&myhandle);
#else
cublasHandle_t myhandle=NULL;
#endif
magma_int_t **dipiv_displ = NULL;
double **dA_displ = NULL;
double **dW0_displ = NULL;
double **dW1_displ = NULL;
double **dW2_displ = NULL;
double **dW3_displ = NULL;
double **dW4_displ = NULL;
double **dinvA_array = NULL;
double **dwork_array = NULL;
magma_malloc((void**)&dipiv_displ, batchCount * sizeof(*dipiv_displ));
magma_malloc((void**)&dA_displ, batchCount * sizeof(*dA_displ));
magma_malloc((void**)&dW0_displ, batchCount * sizeof(*dW0_displ));
magma_malloc((void**)&dW1_displ, batchCount * sizeof(*dW1_displ));
magma_malloc((void**)&dW2_displ, batchCount * sizeof(*dW2_displ));
magma_malloc((void**)&dW3_displ, batchCount * sizeof(*dW3_displ));
magma_malloc((void**)&dW4_displ, batchCount * sizeof(*dW4_displ));
magma_malloc((void**)&dinvA_array, batchCount * sizeof(*dinvA_array));
magma_malloc((void**)&dwork_array, batchCount * sizeof(*dwork_array));
magma_int_t invA_msize = ((n+TRI_NB-1)/TRI_NB)*TRI_NB*TRI_NB;
magma_int_t dwork_msize = n*nb;
magma_int_t **pivinfo_array = NULL;
magma_int_t *pivinfo = NULL;
double* dinvA = NULL;
double* dwork = NULL;// dinvA and dwork are workspace in dtrsm
double **cpuAarray = NULL;
magma_dmalloc( &dinvA, invA_msize * batchCount);
magma_dmalloc( &dwork, dwork_msize * batchCount );
magma_malloc((void**)&pivinfo_array, batchCount * sizeof(*pivinfo_array));
magma_malloc((void**)&pivinfo, batchCount * m * sizeof(magma_int_t));
magma_malloc_cpu((void**) &cpuAarray, batchCount*sizeof(double*));
/* check allocation */
if ( dA_displ == NULL || dW0_displ == NULL || dW1_displ == NULL || dW2_displ == NULL ||
dW3_displ == NULL || dW4_displ == NULL || dinvA_array == NULL || dwork_array == NULL ||
dinvA == NULL || dwork == NULL || cpuAarray == NULL ||
dipiv_displ == NULL || pivinfo_array == NULL || pivinfo == NULL) {
magma_free(dA_displ);
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dinvA_array);
magma_free(dwork_array);
magma_free( dinvA );
magma_free( dwork );
free(cpuAarray);
magma_free(dipiv_displ);
magma_free(pivinfo_array);
magma_free(pivinfo);
magma_int_t info = MAGMA_ERR_DEVICE_ALLOC;
magma_xerbla( __func__, -(info) );
return info;
}
magmablas_dlaset_q(MagmaFull, invA_msize, batchCount, MAGMA_D_ZERO, MAGMA_D_ZERO, dinvA, invA_msize, queue);
magmablas_dlaset_q(MagmaFull, dwork_msize, batchCount, MAGMA_D_ZERO, MAGMA_D_ZERO, dwork, dwork_msize, queue);
dset_pointer(dwork_array, dwork, 1, 0, 0, dwork_msize, batchCount, queue);
dset_pointer(dinvA_array, dinvA, TRI_NB, 0, 0, invA_msize, batchCount, queue);
set_ipointer(pivinfo_array, pivinfo, 1, 0, 0, m, batchCount, queue);
// printf(" I am in dgetrfbatched\n");
magma_queue_t cstream;
magmablasGetKernelStream(&cstream);
magma_int_t streamid;
const magma_int_t nbstreams=32;
magma_queue_t stream[nbstreams];
for(i=0; i<nbstreams; i++){
magma_queue_create( &stream[i] );
}
magma_getvector( batchCount, sizeof(double*), dA_array, 1, cpuAarray, 1);
#if defined(ENABLE_TIMER3)
printf(" I am after malloc\n");
talloc = magma_sync_wtime(0) - talloc;
tloop = magma_sync_wtime(0);
#endif
for(i = 0; i < min_mn; i+=nb)
{
magmablasSetKernelStream(NULL);
ib = min(nb, min_mn-i);
pm = m-i;
magma_idisplace_pointers(dipiv_displ, ipiv_array, ldda, i, 0, batchCount, queue);
magma_ddisplace_pointers(dA_displ, dA_array, ldda, i, i, batchCount, queue);
//===============================================
// panel factorization
//===============================================
#if 0
arginfo = magma_dgetf2_batched(
pm, ib,
dA_displ, ldda,
dW1_displ, dW2_displ, dW3_displ,
dipiv_displ,
info_array, i, batchCount, myhandle);
#else
arginfo = magma_dgetrf_recpanel_batched(
pm, ib, 16,
dA_displ, ldda,
dipiv_displ, pivinfo_array,
dwork_array, nb,
dinvA_array, invA_msize,
dW0_displ, dW1_displ, dW2_displ,
dW3_displ, dW4_displ,
info_array, i,
batchCount, myhandle, queue);
#endif
if(arginfo != 0 ) goto fin;
//===============================================
// end of panel
//===============================================
#define RUN_ALL
#ifdef RUN_ALL
// setup pivinfo before adjusting ipiv
setup_pivinfo_batched(pivinfo_array, dipiv_displ, pm, ib, batchCount, queue);
adjust_ipiv_batched(dipiv_displ, ib, i, batchCount, queue);
// stepinit_ipiv(pivinfo_array, pm, batchCount);// for debug and check swap, it create an ipiv
#if 0
dlaswp_batched( i, dA_displ, ldda,
i, i+ib,
dipiv_displ, pivinfo_array, batchCount);
#else
magma_ddisplace_pointers(dA_displ, dA_array, ldda, i, 0, batchCount, queue);
magma_ddisplace_pointers(dW0_displ, dA_array, ldda, i, 0, batchCount, queue);
magma_dlaswp_rowparallel_batched( i, dA_displ, ldda,
dW0_displ, ldda,
i, i+ib,
pivinfo_array, batchCount, queue);
#endif
if( (i + ib) < n)
{
// swap right side and trsm
magma_ddisplace_pointers(dA_displ, dA_array, ldda, i, i+ib, batchCount, queue);
dset_pointer(dwork_array, dwork, nb, 0, 0, dwork_msize, batchCount, queue); // I don't think it is needed Azzam
magma_dlaswp_rowparallel_batched( n-(i+ib), dA_displ, ldda,
dwork_array, nb,
i, i+ib,
pivinfo_array, batchCount, queue);
magma_ddisplace_pointers(dA_displ, dA_array, ldda, i, i, batchCount, queue);
magma_ddisplace_pointers(dW0_displ, dA_array, ldda, i, i+ib, batchCount, queue);
magmablas_dtrsm_outofplace_batched(MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit, 1,
ib, n-i-ib,
MAGMA_D_ONE,
dA_displ, ldda, // dA
dwork_array, nb, // dB
dW0_displ, ldda, // dX
dinvA_array, invA_msize,
dW1_displ, dW2_displ,
dW3_displ, dW4_displ,
0, batchCount, queue);
if( (i + ib) < m)
{
// if gemm size is >160 use a streamed classical cublas gemm since it is faster
// the batched is faster only when M=N<=160 for K40c
//-------------------------------------------
// USE STREAM GEMM
//-------------------------------------------
if( (m-i-ib) > gemm_crossover && (n-i-ib) > gemm_crossover)
{
//printf("caling streamed dgemm %d %d %d \n", m-i-ib, n-i-ib, ib);
// since it use different stream I need to wait the TRSM and swap.
// But since the code use the NULL stream everywhere,
// so I don't need it, because the NULL stream do the sync by itself
//magma_queue_sync(NULL);
//
for(k=0; k<batchCount; k++)
{
streamid = k%nbstreams;
magmablasSetKernelStream(stream[streamid]);
magma_dgemm(MagmaNoTrans, MagmaNoTrans,
m-i-ib, n-i-ib, ib,
neg_one, cpuAarray[k] + (i+ib)+i*ldda, ldda,
cpuAarray[k] + i+(i+ib)*ldda, ldda,
one, cpuAarray[k] + (i+ib)+(i+ib)*ldda, ldda);
}
// need to synchronise to be sure that dgetf2 do not start before
// finishing the update at least of the next panel
// BUT no need for it as soon as the other portion of the code
// use the NULL stream which do the sync by itself
//magma_device_sync();
}
//-------------------------------------------
// USE BATCHED GEMM
//-------------------------------------------
else
{
magma_ddisplace_pointers(dA_displ, dA_array, ldda, i+ib, i, batchCount, queue);
magma_ddisplace_pointers(dW1_displ, dA_array, ldda, i, i+ib, batchCount, queue);
magma_ddisplace_pointers(dW2_displ, dA_array, ldda, i+ib, i+ib, batchCount, queue);
//printf("caling batched dgemm %d %d %d \n", m-i-ib, n-i-ib, ib);
magmablas_dgemm_batched( MagmaNoTrans, MagmaNoTrans, m-i-ib, n-i-ib, ib,
neg_one, dA_displ, ldda,
dW1_displ, ldda,
one, dW2_displ, ldda,
batchCount, queue);
} // end of batched/stream gemm
} // end of if( (i + ib) < m)
} // end of if( (i + ib) < n)
#endif
}// end of for
fin:
magma_queue_sync(NULL);
#if defined(ENABLE_TIMER3)
tloop = magma_sync_wtime(0) - tloop;
tdalloc = magma_sync_wtime(0);
#endif
for(i=0; i<nbstreams; i++){
magma_queue_destroy( stream[i] );
}
magmablasSetKernelStream(cstream);
#if defined(USE_CUOPT)
cublasDestroy_v2(myhandle);
#endif
magma_free(dA_displ);
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dinvA_array);
magma_free(dwork_array);
magma_free( dinvA );
magma_free( dwork );
free(cpuAarray);
magma_free(dipiv_displ);
magma_free(pivinfo_array);
magma_free(pivinfo);
#if defined(ENABLE_TIMER3)
tdalloc = magma_sync_wtime(0) - tdalloc;
tall = magma_sync_wtime(0) - tall;
printf("here is the timing from inside dgetrf_batched talloc: %10.5f tloop: %10.5f tdalloc: %10.5f tall: %10.5f sum: %10.5f\n", talloc, tloop, tdalloc, tall, talloc+tloop+tdalloc );
#endif
return arginfo;
}
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, t1, t2;
double c_neg_one = MAGMA_D_NEG_ONE;
magma_int_t ione = 1;
const char trans[] = { 'N', 'C', 'T' };
const char uplo[] = { 'L', 'U' };
const char diag[] = { 'U', 'N' };
const char side[] = { 'L', 'R' };
double *A, *B, *C, *C2, *LU;
double *dA, *dB, *dC1, *dC2;
double alpha = MAGMA_D_MAKE( 0.5, 0.1 );
double beta = MAGMA_D_MAKE( 0.7, 0.2 );
double dalpha = 0.6;
double dbeta = 0.8;
double work[1], error, total_error;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t m, n, k, size, maxn, ld, info;
magma_int_t *piv;
magma_err_t err;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf( "Compares magma wrapper function to cublas function; all diffs should be exactly 0.\n\n" );
total_error = 0.;
for( int i = 0; i < opts.ntest; ++i ) {
m = opts.msize[i];
n = opts.nsize[i];
k = opts.ksize[i];
printf("=========================================================================\n");
printf( "M %d, N %d, K %d\n", (int) m, (int) n, (int) k );
// allocate matrices
// over-allocate so they can be any combination of {m,n,k} x {m,n,k}.
maxn = max( max( m, n ), k );
ld = maxn;
size = maxn*maxn;
err = magma_malloc_cpu( (void**) &piv, maxn*sizeof(magma_int_t) ); assert( err == 0 );
err = magma_dmalloc_pinned( &A, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &B, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &C, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &C2, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &LU, size ); assert( err == 0 );
err = magma_dmalloc( &dA, size ); assert( err == 0 );
err = magma_dmalloc( &dB, size ); assert( err == 0 );
err = magma_dmalloc( &dC1, size ); assert( err == 0 );
err = magma_dmalloc( &dC2, size ); assert( err == 0 );
// initialize matrices
size = maxn*maxn;
lapackf77_dlarnv( &ione, ISEED, &size, A );
lapackf77_dlarnv( &ione, ISEED, &size, B );
lapackf77_dlarnv( &ione, ISEED, &size, C );
printf( "========== Level 1 BLAS ==========\n" );
// ----- test DSWAP
// swap 2nd and 3rd columns of dA, then copy to C2 and compare with A
assert( n >= 4 );
magma_dsetmatrix( m, n, A, ld, dA, ld );
magma_dsetmatrix( m, n, A, ld, dB, ld );
magma_dswap( m, dA(0,1), 1, dA(0,2), 1 );
magma_dswap( m, dB(0,1), 1, dB(0,2), 1 );
// check results, storing diff between magma and cuda calls in C2
cublasDaxpy( ld*n, c_neg_one, dA, 1, dB, 1 );
magma_dgetmatrix( m, n, dB, ld, C2, ld );
error = lapackf77_dlange( "F", &m, &k, C2, &ld, work );
total_error += error;
printf( "dswap diff %.2g\n", error );
// ----- test IDAMAX
// get argmax of column of A
magma_dsetmatrix( m, k, A, ld, dA, ld );
error = 0;
for( int j = 0; j < k; ++j ) {
magma_int_t i1 = magma_idamax( m, dA(0,j), 1 );
magma_int_t i2 = cublasIdamax( m, dA(0,j), 1 );
assert( i1 == i2 );
error += abs( i1 - i2 );
}
total_error += error;
gflops = (double)m * k / 1e9;
printf( "idamax diff %.2g\n", error );
printf( "\n" );
printf( "========== Level 2 BLAS ==========\n" );
// ----- test DGEMV
// c = alpha*A*b + beta*c, with A m*n; b,c m or n-vectors
// try no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
magma_dsetmatrix( m, n, A, ld, dA, ld );
magma_dsetvector( maxn, B, 1, dB, 1 );
magma_dsetvector( maxn, C, 1, dC1, 1 );
magma_dsetvector( maxn, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_dgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
size = (trans[ia] == 'N' ? m : n);
cublasDaxpy( size, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetvector( size, dC2, 1, C2, 1 );
error = lapackf77_dlange( "F", &size, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_DGEMV( m, n ) / 1e9;
printf( "dgemv( %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
trans[ia], error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test DSYMV
// c = alpha*A*b + beta*c, with A m*m symmetric; b,c m-vectors
// try upper/lower
for( int iu = 0; iu < 2; ++iu ) {
magma_dsetmatrix( m, m, A, ld, dA, ld );
magma_dsetvector( m, B, 1, dB, 1 );
magma_dsetvector( m, C, 1, dC1, 1 );
magma_dsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_dsymv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDsymv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( m, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_dlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYMV( m ) / 1e9;
printf( "dsymv( %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test DTRSV
// solve A*c = c, with A m*m triangular; c m-vector
// try upper/lower, no-trans/trans, unit/non-unit diag
// Factor A into LU to get well-conditioned triangles, else solve yields garbage.
// Still can give garbage if solves aren't consistent with LU factors,
// e.g., using unit diag for U, so copy lower triangle to upper triangle.
// Also used for trsm later.
lapackf77_dlacpy( "Full", &maxn, &maxn, A, &ld, LU, &ld );
lapackf77_dgetrf( &maxn, &maxn, LU, &ld, piv, &info );
for( int j = 0; j < maxn; ++j ) {
for( int i = 0; i < j; ++i ) {
*LU(i,j) = *LU(j,i);
}
}
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
magma_dsetmatrix( m, m, LU, ld, dA, ld );
magma_dsetvector( m, C, 1, dC1, 1 );
magma_dsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_dtrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDtrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( m, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_dlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_DTRSM( MagmaLeft, m, 1 ) / 1e9;
printf( "dtrsv( %c, %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], diag[id], error, gflops/t1, gflops/t2 );
}}}
printf( "\n" );
printf( "========== Level 3 BLAS ==========\n" );
// ----- test DGEMM
// C = alpha*A*B + beta*C, with A m*k or k*m; B k*n or n*k; C m*n
// try combinations of no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
for( int ib = 0; ib < 3; ++ib ) {
bool nta = (trans[ia] == 'N');
bool ntb = (trans[ib] == 'N');
magma_dsetmatrix( (nta ? m : k), (nta ? m : k), A, ld, dA, ld );
magma_dsetmatrix( (ntb ? k : n), (ntb ? n : k), B, ld, dB, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dgemm( trans[ia], trans[ib], m, n, k, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDgemm( trans[ia], trans[ib], m, n, k, alpha, dA, ld, dB, ld, beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DGEMM( m, n, k ) / 1e9;
printf( "dgemm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
trans[ia], trans[ib], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DSYMM
// C = alpha*A*B + beta*C (left) with A m*m symmetric; B,C m*n; or
// C = alpha*B*A + beta*C (right) with A n*n symmetric; B,C m*n
// try left/right, upper/lower
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
magma_dsetmatrix( m, m, A, ld, dA, ld );
magma_dsetmatrix( m, n, B, ld, dB, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dsymm( side[is], uplo[iu], m, n, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDsymm( side[is], uplo[iu], m, n, alpha, dA, ld, dB, ld, beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYMM( side[is], m, n ) / 1e9;
printf( "dsymm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
side[is], uplo[iu], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DSYRK
// C = alpha*A*A^H + beta*C (no-trans) with A m*k and C m*m symmetric; or
// C = alpha*A^H*A + beta*C (trans) with A k*m and C m*m symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
magma_dsetmatrix( n, k, A, ld, dA, ld );
magma_dsetmatrix( n, n, C, ld, dC1, ld );
magma_dsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dsyrk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDsyrk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYRK( k, n ) / 1e9;
printf( "dsyrk( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DSYR2K
// C = alpha*A*B^H + ^alpha*B*A^H + beta*C (no-trans) with A,B n*k; C n*n symmetric; or
// C = alpha*A^H*B + ^alpha*B^H*A + beta*C (trans) with A,B k*n; C n*n symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
bool nt = (trans[it] == 'N');
magma_dsetmatrix( (nt ? n : k), (nt ? n : k), A, ld, dA, ld );
magma_dsetmatrix( n, n, C, ld, dC1, ld );
magma_dsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dsyr2k( uplo[iu], trans[it], n, k, alpha, dA, ld, dB, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDsyr2k( uplo[iu], trans[it], n, k, alpha, dA, ld, dB, ld, dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYR2K( k, n ) / 1e9;
printf( "dsyr2k( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DTRMM
// C = alpha*A*C (left) with A m*m triangular; C m*n; or
// C = alpha*C*A (right) with A n*n triangular; C m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == 'L');
magma_dsetmatrix( (left ? m : n), (left ? m : n), A, ld, dA, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dtrmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDtrmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DTRMM( side[is], m, n ) / 1e9;
printf( "dtrmm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// ----- test DTRSM
// solve A*X = alpha*B (left) with A m*m triangular; B m*n; or
// solve X*A = alpha*B (right) with A n*n triangular; B m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == 'L');
magma_dsetmatrix( (left ? m : n), (left ? m : n), LU, ld, dA, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dtrsm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDtrsm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DTRSM( side[is], m, n ) / 1e9;
printf( "dtrsm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// cleanup
magma_free_cpu( piv );
magma_free_pinned( A );
magma_free_pinned( B );
magma_free_pinned( C );
magma_free_pinned( C2 );
magma_free_pinned( LU );
magma_free( dA );
magma_free( dB );
magma_free( dC1 );
magma_free( dC2 );
}
if ( total_error != 0. ) {
printf( "total error %.2g -- ought to be 0 -- some test failed (see above).\n",
total_error );
}
else {
printf( "all tests passed\n" );
}
TESTING_FINALIZE();
return 0;
}
extern "C" magma_int_t
magma_slobpcg( magma_s_sparse_matrix A, magma_s_solver_par *solver_par ) {
#define residualNorms(i,iter) ( residualNorms + (i) + (iter)*n )
#define magmablas_swap(x, y) { pointer = x; x = y; y = pointer; }
#define hresidualNorms(i,iter) (hresidualNorms + (i) + (iter)*n )
#define gramA( m, n) (gramA + (m) + (n)*ldgram)
#define gramB( m, n) (gramB + (m) + (n)*ldgram)
#define gevectors(m, n) (gevectors + (m) + (n)*ldgram)
#define h_gramB( m, n) (h_gramB + (m) + (n)*ldgram)
#define magma_s_bspmv_tuned(m, n, alpha, A, X, beta, AX) { \
magmablas_stranspose( m, n, X, m, blockW, n ); \
magma_s_vector x, ax; \
x.memory_location = Magma_DEV; x.num_rows = m*n; x.nnz = m*n; x.val = blockW; \
ax.memory_location= Magma_DEV; ax.num_rows = m*n; ax.nnz = m*n; ax.val = AX; \
magma_s_spmv(alpha, A, x, beta, ax ); \
magmablas_stranspose( n, m, blockW, n, X, m ); \
}
//**************************************************************
// Memory allocation for the eigenvectors, eigenvalues, and workspace
solver_par->solver = Magma_LOBPCG;
magma_int_t m = A.num_rows;
magma_int_t n =(solver_par->num_eigenvalues);
float *blockX = solver_par->eigenvectors;
float *evalues = solver_par->eigenvalues;
float *dwork, *hwork;
float *blockP, *blockAP, *blockR, *blockAR, *blockAX, *blockW;
float *gramA, *gramB, *gramM;
float *gevectors, *h_gramB;
float *pointer, *origX = blockX;
float *eval_gpu;
magma_int_t lwork = max( 2*n+n*magma_get_dsytrd_nb(n),
1 + 6*3*n + 2* 3*n* 3*n);
magma_smalloc_pinned( &hwork , lwork );
magma_smalloc( &blockAX , m*n );
magma_smalloc( &blockAR , m*n );
magma_smalloc( &blockAP , m*n );
magma_smalloc( &blockR , m*n );
magma_smalloc( &blockP , m*n );
magma_smalloc( &blockW , m*n );
magma_smalloc( &dwork , m*n );
magma_smalloc( &eval_gpu , 3*n );
//**********************************************************+
magma_int_t verbosity = 1;
magma_int_t *iwork, liwork = 15*n+9;
// === Set solver parameters ===
float residualTolerance = solver_par->epsilon;
magma_int_t maxIterations = solver_par->maxiter;
// === Set some constants & defaults ===
float c_one = MAGMA_S_ONE, c_zero = MAGMA_S_ZERO;
float *residualNorms, *condestGhistory, condestG;
float *gevalues;
magma_int_t *activeMask;
// === Check some parameters for possible quick exit ===
solver_par->info = 0;
if (m < 2)
solver_par->info = -1;
else if (n > m)
solver_par->info = -2;
if (solver_par->info != 0) {
magma_xerbla( __func__, -(solver_par->info) );
return solver_par->info;
}
magma_int_t *info = &(solver_par->info); // local info variable;
// === Allocate GPU memory for the residual norms' history ===
magma_smalloc(&residualNorms, (maxIterations+1) * n);
magma_malloc( (void **)&activeMask, (n+1) * sizeof(magma_int_t) );
// === Allocate CPU work space ===
magma_smalloc_cpu(&condestGhistory, maxIterations+1);
magma_smalloc_cpu(&gevalues, 3 * n);
magma_malloc_cpu((void **)&iwork, liwork * sizeof(magma_int_t));
float *hW;
magma_smalloc_pinned(&hW, n*n);
magma_smalloc_pinned(&gevectors, 9*n*n);
magma_smalloc_pinned(&h_gramB , 9*n*n);
// === Allocate GPU workspace ===
magma_smalloc(&gramM, n * n);
magma_smalloc(&gramA, 9 * n * n);
magma_smalloc(&gramB, 9 * n * n);
#if defined(PRECISION_z) || defined(PRECISION_c)
float *rwork;
magma_int_t lrwork = 1 + 5*(3*n) + 2*(3*n)*(3*n);
magma_smalloc_cpu(&rwork, lrwork);
#endif
// === Set activemask to one ===
for(int k =0; k<n; k++)
iwork[k]=1;
magma_setmatrix(n, 1, sizeof(magma_int_t), iwork, n ,activeMask, n);
magma_int_t gramDim, ldgram = 3*n, ikind = 4;
// === Make the initial vectors orthonormal ===
magma_sgegqr_gpu(ikind, m, n, blockX, m, dwork, hwork, info );
//magma_sorthomgs( m, n, blockX );
magma_s_bspmv_tuned(m, n, c_one, A, blockX, c_zero, blockAX );
// === Compute the Gram matrix = (X, AX) & its eigenstates ===
magma_sgemm(MagmaTrans, MagmaNoTrans, n, n, m,
c_one, blockX, m, blockAX, m, c_zero, gramM, n);
magma_ssyevd_gpu( MagmaVec, MagmaUpper,
n, gramM, n, evalues, hW, n, hwork, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork, info );
// === Update X = X * evectors ===
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockX, m, gramM, n, c_zero, blockW, m);
magmablas_swap(blockW, blockX);
// === Update AX = AX * evectors ===
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockAX, m, gramM, n, c_zero, blockW, m);
magmablas_swap(blockW, blockAX);
condestGhistory[1] = 7.82;
magma_int_t iterationNumber, cBlockSize, restart = 1, iter;
//Chronometry
real_Double_t tempo1, tempo2;
magma_device_sync();
tempo1=magma_wtime();
// === Main LOBPCG loop ============================================================
for(iterationNumber = 1; iterationNumber < maxIterations; iterationNumber++)
{
// === compute the residuals (R = Ax - x evalues )
magmablas_slacpy( MagmaUpperLower, m, n, blockAX, m, blockR, m);
/*
for(int i=0; i<n; i++){
magma_saxpy(m, MAGMA_S_MAKE(-evalues[i],0), blockX+i*m, 1, blockR+i*m, 1);
}
*/
#if defined(PRECISION_z) || defined(PRECISION_d)
magma_dsetmatrix( 3*n, 1, evalues, 3*n, eval_gpu, 3*n );
#else
magma_ssetmatrix( 3*n, 1, evalues, 3*n, eval_gpu, 3*n );
#endif
magma_slobpcg_res( m, n, eval_gpu, blockX, blockR, eval_gpu);
magmablas_snrm2_cols(m, n, blockR, m, residualNorms(0, iterationNumber));
// === remove the residuals corresponding to already converged evectors
magma_scompact(m, n, blockR, m,
residualNorms(0, iterationNumber), residualTolerance,
activeMask, &cBlockSize);
if (cBlockSize == 0)
break;
// === apply a preconditioner P to the active residulas: R_new = P R_old
// === for now set P to be identity (no preconditioner => nothing to be done )
// magmablas_slacpy( MagmaUpperLower, m, cBlockSize, blockR, m, blockW, m);
/*
// === make the preconditioned residuals orthogonal to X
magma_sgemm(MagmaTrans, MagmaNoTrans, n, cBlockSize, m,
c_one, blockX, m, blockR, m, c_zero, gramB(0,0), ldgram);
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, cBlockSize, n,
c_mone, blockX, m, gramB(0,0), ldgram, c_one, blockR, m);
*/
// === make the active preconditioned residuals orthonormal
magma_sgegqr_gpu(ikind, m, cBlockSize, blockR, m, dwork, hwork, info );
//magma_sorthomgs( m, cBlockSize, blockR );
// === compute AR
magma_s_bspmv_tuned(m, cBlockSize, c_one, A, blockR, c_zero, blockAR );
if (!restart) {
// === compact P & AP as well
magma_scompactActive(m, n, blockP, m, activeMask);
magma_scompactActive(m, n, blockAP, m, activeMask);
/*
// === make P orthogonal to X ?
magma_sgemm(MagmaTrans, MagmaNoTrans, n, cBlockSize, m,
c_one, blockX, m, blockP, m, c_zero, gramB(0,0), ldgram);
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, cBlockSize, n,
c_mone, blockX, m, gramB(0,0), ldgram, c_one, blockP, m);
// === make P orthogonal to R ?
magma_sgemm(MagmaTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockR, m, blockP, m, c_zero, gramB(0,0), ldgram);
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, cBlockSize, cBlockSize,
c_mone, blockR, m, gramB(0,0), ldgram, c_one, blockP, m);
*/
// === Make P orthonormal & properly change AP (without multiplication by A)
magma_sgegqr_gpu(ikind, m, cBlockSize, blockP, m, dwork, hwork, info );
//magma_sorthomgs( m, cBlockSize, blockP );
//magma_s_bspmv_tuned(m, cBlockSize, c_one, A, blockP, c_zero, blockAP );
magma_ssetmatrix( cBlockSize, cBlockSize, hwork, cBlockSize, dwork, cBlockSize);
// magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
// m, cBlockSize, c_one, dwork, cBlockSize, blockAP, m);
// replacement according to Stan
#if defined(PRECISION_s) || defined(PRECISION_d)
magmablas_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
m, cBlockSize, c_one, dwork, cBlockSize, blockAP, m);
#else
magma_strsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit, m,
cBlockSize, c_one, dwork, cBlockSize, blockAP, m);
#endif
}
iter = max(1,iterationNumber-10- (int)(log(1.*cBlockSize)));
float condestGmean = 0.;
for(int i = 0; i<iterationNumber-iter+1; i++)
condestGmean += condestGhistory[i];
condestGmean = condestGmean / (iterationNumber-iter+1);
if (restart)
gramDim = n+cBlockSize;
else
gramDim = n+2*cBlockSize;
/* --- The Raileight-Ritz method for [X R P] -----------------------
[ X R P ]' [AX AR AP] y = evalues [ X R P ]' [ X R P ], i.e.,
GramA GramB
/ X'AX X'AR X'AP \ / X'X X'R X'P \
| R'AX R'AR R'AP | y = evalues | R'X R'R R'P |
\ P'AX P'AR P'AP / \ P'X P'R P'P /
----------------------------------------------------------------- */
// === assemble GramB; first, set it to I
magmablas_slaset(MagmaFull, ldgram, ldgram, c_zero, c_one, gramB, ldgram); // identity
if (!restart) {
magma_sgemm(MagmaTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockP, m, blockX, m, c_zero, gramB(n+cBlockSize,0), ldgram);
magma_sgemm(MagmaTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockP, m, blockR, m, c_zero, gramB(n+cBlockSize,n), ldgram);
}
magma_sgemm(MagmaTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockR, m, blockX, m, c_zero, gramB(n,0), ldgram);
// === get GramB from the GPU to the CPU and compute its eigenvalues only
magma_sgetmatrix(gramDim, gramDim, gramB, ldgram, h_gramB, ldgram);
lapackf77_ssyev("N", "L", &gramDim, h_gramB, &ldgram, gevalues,
hwork, &lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork,
#endif
info);
// === check stability criteria if we need to restart
condestG = log10( gevalues[gramDim-1]/gevalues[0] ) + 1.;
if ((condestG/condestGmean>2 && condestG>2) || condestG>8) {
// Steepest descent restart for stability
restart=1;
printf("restart at step #%d\n", (int) iterationNumber);
}
// === assemble GramA; first, set it to I
magmablas_slaset(MagmaFull, ldgram, ldgram, c_zero, c_one, gramA, ldgram); // identity
magma_sgemm(MagmaTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockR, m, blockAX, m, c_zero, gramA(n,0), ldgram);
magma_sgemm(MagmaTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockR, m, blockAR, m, c_zero, gramA(n,n), ldgram);
if (!restart) {
magma_sgemm(MagmaTrans, MagmaNoTrans, cBlockSize, n, m,
c_one, blockP, m, blockAX, m, c_zero,
gramA(n+cBlockSize,0), ldgram);
magma_sgemm(MagmaTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockP, m, blockAR, m, c_zero,
gramA(n+cBlockSize,n), ldgram);
magma_sgemm(MagmaTrans, MagmaNoTrans, cBlockSize, cBlockSize, m,
c_one, blockP, m, blockAP, m, c_zero,
gramA(n+cBlockSize,n+cBlockSize), ldgram);
}
/*
// === Compute X' AX or just use the eigenvalues below ?
magma_sgemm(MagmaTrans, MagmaNoTrans, n, n, m,
c_one, blockX, m, blockAX, m, c_zero,
gramA(0,0), ldgram);
*/
if (restart==0) {
magma_sgetmatrix(gramDim, gramDim, gramA, ldgram, gevectors, ldgram);
}
else {
gramDim = n+cBlockSize;
magma_sgetmatrix(gramDim, gramDim, gramA, ldgram, gevectors, ldgram);
}
for(int k=0; k<n; k++)
*gevectors(k,k) = MAGMA_S_MAKE(evalues[k], 0);
// === the previous eigensolver destroyed what is in h_gramB => must copy it again
magma_sgetmatrix(gramDim, gramDim, gramB, ldgram, h_gramB, ldgram);
magma_int_t itype = 1;
lapackf77_ssygvd(&itype, "V", "L", &gramDim,
gevectors, &ldgram, h_gramB, &ldgram,
gevalues, hwork, &lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, &lrwork,
#endif
iwork, &liwork, info);
for(int k =0; k<n; k++)
evalues[k] = gevalues[k];
// === copy back the result to gramA on the GPU and use it for the updates
magma_ssetmatrix(gramDim, gramDim, gevectors, ldgram, gramA, ldgram);
if (restart == 0) {
// === contribution from P to the new X (in new search direction P)
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockP, m, gramA(n+cBlockSize,0), ldgram, c_zero, dwork, m);
magmablas_swap(dwork, blockP);
// === contribution from R to the new X (in new search direction P)
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockR, m, gramA(n,0), ldgram, c_one, blockP, m);
// === corresponding contribution from AP to the new AX (in AP)
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockAP, m, gramA(n+cBlockSize,0), ldgram, c_zero, dwork, m);
magmablas_swap(dwork, blockAP);
// === corresponding contribution from AR to the new AX (in AP)
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockAR, m, gramA(n,0), ldgram, c_one, blockAP, m);
}
else {
// === contribution from R (only) to the new X
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, cBlockSize,
c_one, blockR, m, gramA(n,0), ldgram, c_zero, blockP, m);
// === corresponding contribution from AR (only) to the new AX
magma_sgemm(MagmaNoTrans, MagmaNoTrans,m, n, cBlockSize,
c_one, blockAR, m, gramA(n,0), ldgram, c_zero, blockAP, m);
}
// === contribution from old X to the new X + the new search direction P
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockX, m, gramA, ldgram, c_zero, dwork, m);
magmablas_swap(dwork, blockX);
//magma_saxpy(m*n, c_one, blockP, 1, blockX, 1);
magma_slobpcg_maxpy( m, n, blockP, blockX );
// === corresponding contribution from old AX to new AX + AP
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, blockAX, m, gramA, ldgram, c_zero, dwork, m);
magmablas_swap(dwork, blockAX);
//magma_saxpy(m*n, c_one, blockAP, 1, blockAX, 1);
magma_slobpcg_maxpy( m, n, blockAP, blockAX );
condestGhistory[iterationNumber+1]=condestG;
if (verbosity==1) {
// float res;
// magma_sgetmatrix(1, 1,
// (float*)residualNorms(0, iterationNumber), 1,
// (float*)&res, 1);
//
// printf("Iteration %4d, CBS %4d, Residual: %10.7f\n",
// iterationNumber, cBlockSize, res);
printf("%4d-%2d ", (int) iterationNumber, (int) cBlockSize);
magma_sprint_gpu(1, n, residualNorms(0, iterationNumber), 1);
}
restart = 0;
} // === end for iterationNumber = 1,maxIterations =======================
// fill solver info
magma_device_sync();
tempo2=magma_wtime();
solver_par->runtime = (real_Double_t) tempo2-tempo1;
solver_par->numiter = iterationNumber;
if( solver_par->numiter < solver_par->maxiter) {
solver_par->info = 0;
} else if( solver_par->init_res > solver_par->final_res )
solver_par->info = -2;
else
solver_par->info = -1;
// =============================================================================
// === postprocessing;
// =============================================================================
// === compute the real AX and corresponding eigenvalues
magma_s_bspmv_tuned(m, n, c_one, A, blockX, c_zero, blockAX );
magma_sgemm(MagmaTrans, MagmaNoTrans, n, n, m,
c_one, blockX, m, blockAX, m, c_zero, gramM, n);
magma_ssyevd_gpu( MagmaVec, MagmaUpper,
n, gramM, n, gevalues, dwork, n, hwork, lwork,
#if defined(PRECISION_z) || defined(PRECISION_c)
rwork, lrwork,
#endif
iwork, liwork, info );
for(int k =0; k<n; k++)
evalues[k] = gevalues[k];
// === update X = X * evectors
magmablas_swap(blockX, dwork);
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, dwork, m, gramM, n, c_zero, blockX, m);
// === update AX = AX * evectors to compute the final residual
magmablas_swap(blockAX, dwork);
magma_sgemm(MagmaNoTrans, MagmaNoTrans, m, n, n,
c_one, dwork, m, gramM, n, c_zero, blockAX, m);
// === compute R = AX - evalues X
magmablas_slacpy( MagmaUpperLower, m, n, blockAX, m, blockR, m);
for(int i=0; i<n; i++)
magma_saxpy(m, MAGMA_S_MAKE(-evalues[i], 0), blockX+i*m, 1, blockR+i*m, 1);
// === residualNorms[iterationNumber] = || R ||
magmablas_snrm2_cols(m, n, blockR, m, residualNorms(0, iterationNumber));
// === restore blockX if needed
if (blockX != origX)
magmablas_slacpy( MagmaUpperLower, m, n, blockX, m, origX, m);
printf("Eigenvalues:\n");
for(int i =0; i<n; i++)
printf("%e ", evalues[i]);
printf("\n\n");
printf("Final residuals:\n");
magma_sprint_gpu(1, n, residualNorms(0, iterationNumber), 1);
printf("\n\n");
//=== Print residual history in a file for plotting ====
float *hresidualNorms;
magma_smalloc_cpu(&hresidualNorms, (iterationNumber+1) * n);
magma_sgetmatrix(n, iterationNumber,
(float*)residualNorms, n,
(float*)hresidualNorms, n);
printf("Residuals are stored in file residualNorms\n");
printf("Plot the residuals using: myplot \n");
FILE *residuals_file;
residuals_file = fopen("residualNorms", "w");
for(int i =1; i<iterationNumber; i++) {
for(int j = 0; j<n; j++)
fprintf(residuals_file, "%f ", *hresidualNorms(j,i));
fprintf(residuals_file, "\n");
}
fclose(residuals_file);
magma_free_cpu(hresidualNorms);
// === free work space
magma_free( residualNorms );
magma_free_cpu( condestGhistory );
magma_free_cpu( gevalues );
magma_free_cpu( iwork );
magma_free_pinned( hW );
magma_free_pinned( gevectors );
magma_free_pinned( h_gramB );
magma_free( gramM );
magma_free( gramA );
magma_free( gramB );
magma_free( activeMask );
magma_free( blockAX );
magma_free( blockAR );
magma_free( blockAP );
magma_free( blockR );
magma_free( blockP );
magma_free( blockW );
magma_free( dwork );
magma_free( eval_gpu );
magma_free_pinned( hwork );
#if defined(PRECISION_z) || defined(PRECISION_c)
magma_free_cpu( rwork );
#endif
return MAGMA_SUCCESS;
}
/* ////////////////////////////////////////////////////////////////////////////
-- testing sparse matrix vector product
*/
int main( int argc, char** argv )
{
magma_int_t info = 0;
TESTING_CHECK( magma_init() );
magma_print_environment();
magma_queue_t queue=NULL;
magma_queue_create( 0, &queue );
magma_s_matrix hA={Magma_CSR}, hA_SELLP={Magma_CSR},
dA={Magma_CSR}, dA_SELLP={Magma_CSR};
magma_s_matrix hx={Magma_CSR}, hy={Magma_CSR}, dx={Magma_CSR},
dy={Magma_CSR}, hrefvec={Magma_CSR}, hcheck={Magma_CSR};
hA_SELLP.blocksize = 8;
hA_SELLP.alignment = 8;
real_Double_t start, end, res;
#ifdef MAGMA_WITH_MKL
magma_int_t *pntre=NULL;
#endif
cusparseHandle_t cusparseHandle = NULL;
cusparseMatDescr_t descr = NULL;
float c_one = MAGMA_S_MAKE(1.0, 0.0);
float c_zero = MAGMA_S_MAKE(0.0, 0.0);
float accuracy = 1e-10;
#define PRECISION_s
#if defined(PRECISION_c)
accuracy = 1e-4;
#endif
#if defined(PRECISION_s)
accuracy = 1e-4;
#endif
magma_int_t i, j;
for( i = 1; i < argc; ++i ) {
if ( strcmp("--blocksize", argv[i]) == 0 ) {
hA_SELLP.blocksize = atoi( argv[++i] );
} else if ( strcmp("--alignment", argv[i]) == 0 ) {
hA_SELLP.alignment = atoi( argv[++i] );
} else
break;
}
printf("\n# usage: ./run_sspmm"
" [ --blocksize %lld --alignment %lld (for SELLP) ] matrices\n\n",
(long long) hA_SELLP.blocksize, (long long) hA_SELLP.alignment );
while( i < argc ) {
if ( strcmp("LAPLACE2D", argv[i]) == 0 && i+1 < argc ) { // Laplace test
i++;
magma_int_t laplace_size = atoi( argv[i] );
TESTING_CHECK( magma_sm_5stencil( laplace_size, &hA, queue ));
} else { // file-matrix test
TESTING_CHECK( magma_s_csr_mtx( &hA, argv[i], queue ));
}
printf("%% matrix info: %lld-by-%lld with %lld nonzeros\n",
(long long) hA.num_rows, (long long) hA.num_cols, (long long) hA.nnz );
real_Double_t FLOPS = 2.0*hA.nnz/1e9;
// m - number of rows for the sparse matrix
// n - number of vectors to be multiplied in the SpMM product
magma_int_t m, n;
m = hA.num_rows;
n = 48;
// init CPU vectors
TESTING_CHECK( magma_svinit( &hx, Magma_CPU, m, n, c_one, queue ));
TESTING_CHECK( magma_svinit( &hy, Magma_CPU, m, n, c_zero, queue ));
// init DEV vectors
TESTING_CHECK( magma_svinit( &dx, Magma_DEV, m, n, c_one, queue ));
TESTING_CHECK( magma_svinit( &dy, Magma_DEV, m, n, c_zero, queue ));
// calling MKL with CSR
#ifdef MAGMA_WITH_MKL
TESTING_CHECK( magma_imalloc_cpu( &pntre, m + 1 ) );
pntre[0] = 0;
for (j=0; j < m; j++ ) {
pntre[j] = hA.row[j+1];
}
MKL_INT num_rows = hA.num_rows;
MKL_INT num_cols = hA.num_cols;
MKL_INT nnz = hA.nnz;
MKL_INT num_vecs = n;
MKL_INT *col;
TESTING_CHECK( magma_malloc_cpu( (void**) &col, nnz * sizeof(MKL_INT) ));
for( magma_int_t t=0; t < hA.nnz; ++t ) {
col[ t ] = hA.col[ t ];
}
MKL_INT *row;
TESTING_CHECK( magma_malloc_cpu( (void**) &row, num_rows * sizeof(MKL_INT) ));
for( magma_int_t t=0; t < hA.num_rows; ++t ) {
row[ t ] = hA.col[ t ];
}
// === Call MKL with consecutive SpMVs, using mkl_scsrmv ===
// warmp up
mkl_scsrmv( "N", &num_rows, &num_cols,
MKL_ADDR(&c_one), "GFNC", MKL_ADDR(hA.val), col, row, pntre,
MKL_ADDR(hx.val),
MKL_ADDR(&c_zero), MKL_ADDR(hy.val) );
start = magma_wtime();
for (j=0; j < 10; j++ ) {
mkl_scsrmv( "N", &num_rows, &num_cols,
MKL_ADDR(&c_one), "GFNC", MKL_ADDR(hA.val), col, row, pntre,
MKL_ADDR(hx.val),
MKL_ADDR(&c_zero), MKL_ADDR(hy.val) );
}
end = magma_wtime();
printf( "\n > MKL SpMVs : %.2e seconds %.2e GFLOP/s (CSR).\n",
(end-start)/10, FLOPS*10/(end-start) );
// === Call MKL with blocked SpMVs, using mkl_scsrmm ===
char transa = 'n';
MKL_INT ldb = n, ldc=n;
char matdescra[6] = {'g', 'l', 'n', 'c', 'x', 'x'};
// warm up
mkl_scsrmm( &transa, &num_rows, &num_vecs, &num_cols, MKL_ADDR(&c_one), matdescra,
MKL_ADDR(hA.val), col, row, pntre,
MKL_ADDR(hx.val), &ldb,
MKL_ADDR(&c_zero),
MKL_ADDR(hy.val), &ldc );
start = magma_wtime();
for (j=0; j < 10; j++ ) {
mkl_scsrmm( &transa, &num_rows, &num_vecs, &num_cols, MKL_ADDR(&c_one), matdescra,
MKL_ADDR(hA.val), col, row, pntre,
MKL_ADDR(hx.val), &ldb,
MKL_ADDR(&c_zero),
MKL_ADDR(hy.val), &ldc );
}
end = magma_wtime();
printf( "\n > MKL SpMM : %.2e seconds %.2e GFLOP/s (CSR).\n",
(end-start)/10, FLOPS*10.*n/(end-start) );
magma_free_cpu( row );
magma_free_cpu( col );
row = NULL;
col = NULL;
#endif // MAGMA_WITH_MKL
// copy matrix to GPU
TESTING_CHECK( magma_smtransfer( hA, &dA, Magma_CPU, Magma_DEV, queue ));
// SpMV on GPU (CSR)
start = magma_sync_wtime( queue );
for (j=0; j < 10; j++) {
TESTING_CHECK( magma_s_spmv( c_one, dA, dx, c_zero, dy, queue ));
}
end = magma_sync_wtime( queue );
printf( " > MAGMA: %.2e seconds %.2e GFLOP/s (standard CSR).\n",
(end-start)/10, FLOPS*10.*n/(end-start) );
TESTING_CHECK( magma_smtransfer( dy, &hrefvec , Magma_DEV, Magma_CPU, queue ));
magma_smfree(&dA, queue );
// convert to SELLP and copy to GPU
TESTING_CHECK( magma_smconvert( hA, &hA_SELLP, Magma_CSR, Magma_SELLP, queue ));
TESTING_CHECK( magma_smtransfer( hA_SELLP, &dA_SELLP, Magma_CPU, Magma_DEV, queue ));
magma_smfree(&hA_SELLP, queue );
magma_smfree( &dy, queue );
TESTING_CHECK( magma_svinit( &dy, Magma_DEV, dx.num_rows, dx.num_cols, c_zero, queue ));
// SpMV on GPU (SELLP)
start = magma_sync_wtime( queue );
for (j=0; j < 10; j++) {
TESTING_CHECK( magma_s_spmv( c_one, dA_SELLP, dx, c_zero, dy, queue ));
}
end = magma_sync_wtime( queue );
printf( " > MAGMA: %.2e seconds %.2e GFLOP/s (SELLP).\n",
(end-start)/10, FLOPS*10.*n/(end-start) );
TESTING_CHECK( magma_smtransfer( dy, &hcheck , Magma_DEV, Magma_CPU, queue ));
res = 0.0;
for(magma_int_t k=0; k < hA.num_rows; k++ ) {
res=res + MAGMA_S_REAL(hcheck.val[k]) - MAGMA_S_REAL(hrefvec.val[k]);
}
printf("%% |x-y|_F = %8.2e\n", res);
if ( res < accuracy )
printf("%% tester spmm SELL-P: ok\n");
else
printf("%% tester spmm SELL-P: failed\n");
magma_smfree( &hcheck, queue );
magma_smfree(&dA_SELLP, queue );
// SpMV on GPU (CUSPARSE - CSR)
// CUSPARSE context //
magma_smfree( &dy, queue );
TESTING_CHECK( magma_svinit( &dy, Magma_DEV, dx.num_rows, dx.num_cols, c_zero, queue ));
//#ifdef PRECISION_d
start = magma_sync_wtime( queue );
TESTING_CHECK( cusparseCreate( &cusparseHandle ));
TESTING_CHECK( cusparseSetStream( cusparseHandle, magma_queue_get_cuda_stream(queue) ));
TESTING_CHECK( cusparseCreateMatDescr( &descr ));
TESTING_CHECK( cusparseSetMatType( descr, CUSPARSE_MATRIX_TYPE_GENERAL ));
TESTING_CHECK( cusparseSetMatIndexBase( descr, CUSPARSE_INDEX_BASE_ZERO ));
float alpha = c_one;
float beta = c_zero;
// copy matrix to GPU
TESTING_CHECK( magma_smtransfer( hA, &dA, Magma_CPU, Magma_DEV, queue) );
for (j=0; j < 10; j++) {
cusparseScsrmm(cusparseHandle,
CUSPARSE_OPERATION_NON_TRANSPOSE,
dA.num_rows, n, dA.num_cols, dA.nnz,
&alpha, descr, dA.dval, dA.drow, dA.dcol,
dx.dval, dA.num_cols, &beta, dy.dval, dA.num_cols);
}
end = magma_sync_wtime( queue );
printf( " > CUSPARSE: %.2e seconds %.2e GFLOP/s (CSR).\n",
(end-start)/10, FLOPS*10*n/(end-start) );
TESTING_CHECK( magma_smtransfer( dy, &hcheck , Magma_DEV, Magma_CPU, queue ));
res = 0.0;
for(magma_int_t k=0; k < hA.num_rows; k++ ) {
res = res + MAGMA_S_REAL(hcheck.val[k]) - MAGMA_S_REAL(hrefvec.val[k]);
}
printf("%% |x-y|_F = %8.2e\n", res);
if ( res < accuracy )
printf("%% tester spmm cuSPARSE: ok\n");
else
printf("%% tester spmm cuSPARSE: failed\n");
magma_smfree( &hcheck, queue );
cusparseDestroyMatDescr( descr );
cusparseDestroy( cusparseHandle );
descr = NULL;
cusparseHandle = NULL;
//#endif
printf("\n\n");
// free CPU memory
magma_smfree( &hA, queue );
magma_smfree( &hx, queue );
magma_smfree( &hy, queue );
magma_smfree( &hrefvec, queue );
// free GPU memory
magma_smfree( &dx, queue );
magma_smfree( &dy, queue );
magma_smfree( &dA, queue);
#ifdef MAGMA_WITH_MKL
magma_free_cpu( pntre );
#endif
i++;
}
magma_queue_destroy( queue );
TESTING_CHECK( magma_finalize() );
return info;
}
extern "C" magma_int_t
magma_dbulge_back(
magma_uplo_t uplo,
magma_int_t n, magma_int_t nb,
magma_int_t ne, magma_int_t Vblksiz,
double *Z, magma_int_t ldz,
magmaDouble_ptr dZ, magma_int_t lddz,
double *V, magma_int_t ldv,
double *TAU,
double *T, magma_int_t ldt,
magma_int_t* info)
{
magma_int_t threads = magma_get_parallel_numthreads();
magma_int_t mklth = magma_get_lapack_numthreads();
magma_set_lapack_numthreads(1);
real_Double_t timeaplQ2=0.0;
double f= 1.;
magma_int_t n_gpu = ne;
//#if defined(PRECISION_s) || defined(PRECISION_d)
//double gpu_cpu_perf = 50; // gpu over cpu performance //100% ev // SandyB. - Kepler (K20c)
//double gpu_cpu_perf = 16; // gpu over cpu performance //100% ev // SandyB. - Fermi (M2090)
//#else
// double gpu_cpu_perf = 27.5; // gpu over cpu performance //100% ev // Westmere - Fermi (M2090)
//double gpu_cpu_perf = 37; // gpu over cpu performance //100% ev // SandyB. - Kepler (K20c)
// double gpu_cpu_perf = 130; // gpu over cpu performance //100% ev // Bulldozer - Kepler (K20X)
//#endif
magma_int_t gpu_cpu_perf = magma_get_dbulge_gcperf();
if (threads > 1) {
f = 1. / (1. + (double)(threads-1)/ ((double)gpu_cpu_perf) );
n_gpu = (magma_int_t)(f*ne);
}
/****************************************************
* apply V2 from left to the eigenvectors Z. dZ = (I-V2*T2*V2')*Z
* **************************************************/
//$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$
//$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$
//$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$
//n_gpu=ne;
//$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$
//$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$
//$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$
//$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$$
timeaplQ2 = magma_wtime();
/*============================
* use GPU+CPU's
*==========================*/
if (n_gpu < ne) {
// define the size of Q to be done on CPU's and the size on GPU's
// note that GPU use Q(1:N_GPU) and CPU use Q(N_GPU+1:N)
#ifdef ENABLE_DEBUG
printf("---> calling GPU + CPU(if N_CPU > 0) to apply V2 to Z with NE %d N_GPU %d N_CPU %d\n",ne, n_gpu, ne-n_gpu);
#endif
magma_dapplyQ_data data_applyQ;
magma_dapplyQ_data_init(&data_applyQ, threads, n, ne, n_gpu, nb, Vblksiz, Z, ldz, V, ldv, TAU, T, ldt, dZ, lddz);
magma_dapplyQ_id_data* arg;
magma_malloc_cpu((void**) &arg, threads*sizeof(magma_dapplyQ_id_data));
pthread_t* thread_id;
magma_malloc_cpu((void**) &thread_id, threads*sizeof(pthread_t));
pthread_attr_t thread_attr;
// ===============================
// relaunch thread to apply Q
// ===============================
// Set one thread per core
pthread_attr_init(&thread_attr);
pthread_attr_setscope(&thread_attr, PTHREAD_SCOPE_SYSTEM);
pthread_setconcurrency(threads);
// Launch threads
for (magma_int_t thread = 1; thread < threads; thread++) {
magma_dapplyQ_id_data_init(&(arg[thread]), thread, &data_applyQ);
pthread_create(&thread_id[thread], &thread_attr, magma_dapplyQ_parallel_section, &arg[thread]);
}
magma_dapplyQ_id_data_init(&(arg[0]), 0, &data_applyQ);
magma_dapplyQ_parallel_section(&arg[0]);
// Wait for completion
for (magma_int_t thread = 1; thread < threads; thread++) {
void *exitcodep;
pthread_join(thread_id[thread], &exitcodep);
}
magma_free_cpu(thread_id);
magma_free_cpu(arg);
magma_dapplyQ_data_destroy(&data_applyQ);
magma_dsetmatrix(n, ne-n_gpu, Z + n_gpu*ldz, ldz, dZ + n_gpu*ldz, lddz);
/*============================
* use only GPU
*==========================*/
} else {
magma_dsetmatrix(n, ne, Z, ldz, dZ, lddz);
magma_dbulge_applyQ_v2(MagmaLeft, ne, n, nb, Vblksiz, dZ, lddz, V, ldv, T, ldt, info);
magma_device_sync();
}
timeaplQ2 = magma_wtime()-timeaplQ2;
magma_set_lapack_numthreads(mklth);
return MAGMA_SUCCESS;
}
/***************************************************************************//**
Purpose
-------
SPOTRF computes the Cholesky factorization of a real symmetric
positive definite matrix dA.
The factorization has the form
dA = U**H * U, if UPLO = MagmaUpper, or
dA = L * L**H, if UPLO = MagmaLower,
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of dA is stored;
- = MagmaLower: Lower triangle of dA is stored.
@param[in]
n INTEGER
The order of the matrix dA. N >= 0.
@param[in,out]
dA_array Array of pointers, dimension (batchCount).
Each is a REAL array on the GPU, dimension (LDDA,N)
On entry, each pointer is a symmetric matrix dA.
If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of dA contains the upper
triangular part of the matrix dA, and the strictly lower
triangular part of dA is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of dA contains the lower
triangular part of the matrix dA, and the strictly upper
triangular part of dA is not referenced.
\n
On exit, if corresponding entry in info_array = 0,
each pointer is 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 each array dA. LDDA >= max(1,N).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param[out]
info_array Array of INTEGERs, dimension (batchCount), for corresponding matrices.
- = 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.
@param[in]
batchCount INTEGER
The number of matrices to operate on.
@param[in]
queue magma_queue_t
Queue to execute in.
@ingroup magma_potrf_batched
*******************************************************************************/
extern "C" magma_int_t
magma_spotrf_lg_batched(
magma_uplo_t uplo, magma_int_t n,
float **dA_array, magma_int_t ldda,
magma_int_t *info_array, magma_int_t batchCount, magma_queue_t queue)
{
magma_int_t arginfo = 0;
#define A(i_, j_) (A + (i_) + (j_)*ldda)
float d_alpha = -1.0;
float d_beta = 1.0;
if ( n > 2048 ) {
#ifndef MAGMA_NOWARNING
printf("=========================================================================================\n"
" WARNING batched routines are designed for small sizes. It might be better to use the\n"
" Native/Hybrid classical routines if you want good performance.\n"
"=========================================================================================\n");
#endif
}
magma_int_t j, k, ib, use_stream;
magma_int_t nb, recnb;
magma_get_spotrf_batched_nbparam(n, &nb, &recnb);
float **dA_displ = NULL;
float **dW0_displ = NULL;
float **dW1_displ = NULL;
float **dW2_displ = NULL;
float **dW3_displ = NULL;
float **dW4_displ = NULL;
float **dinvA_array = NULL;
float **dwork_array = NULL;
magma_malloc((void**)&dA_displ, batchCount * sizeof(*dA_displ));
magma_malloc((void**)&dW0_displ, batchCount * sizeof(*dW0_displ));
magma_malloc((void**)&dW1_displ, batchCount * sizeof(*dW1_displ));
magma_malloc((void**)&dW2_displ, batchCount * sizeof(*dW2_displ));
magma_malloc((void**)&dW3_displ, batchCount * sizeof(*dW3_displ));
magma_malloc((void**)&dW4_displ, batchCount * sizeof(*dW4_displ));
magma_malloc((void**)&dinvA_array, batchCount * sizeof(*dinvA_array));
magma_malloc((void**)&dwork_array, batchCount * sizeof(*dwork_array));
magma_int_t invA_msize = magma_roundup( n, STRTRI_BATCHED_NB )*STRTRI_BATCHED_NB;
magma_int_t dwork_msize = n*nb;
float* dinvA = NULL;
float* dwork = NULL; // dinvA and dwork are workspace in strsm
float **cpuAarray = NULL;
magma_smalloc( &dinvA, invA_msize * batchCount);
magma_smalloc( &dwork, dwork_msize * batchCount );
magma_malloc_cpu((void**) &cpuAarray, batchCount*sizeof(float*));
/* check allocation */
if ( dA_displ == NULL || dW0_displ == NULL || dW1_displ == NULL || dW2_displ == NULL ||
dW3_displ == NULL || dW4_displ == NULL || dinvA_array == NULL || dwork_array == NULL ||
dinvA == NULL || dwork == NULL || cpuAarray == NULL ) {
magma_free(dA_displ);
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dinvA_array);
magma_free(dwork_array);
magma_free( dinvA );
magma_free( dwork );
magma_free_cpu(cpuAarray);
magma_int_t info = MAGMA_ERR_DEVICE_ALLOC;
magma_xerbla( __func__, -(info) );
return info;
}
magmablas_slaset_q( MagmaFull, invA_msize, batchCount, MAGMA_S_ZERO, MAGMA_S_ZERO, dinvA, invA_msize, queue );
magmablas_slaset_q( MagmaFull, dwork_msize, batchCount, MAGMA_S_ZERO, MAGMA_S_ZERO, dwork, dwork_msize, queue );
magma_sset_pointer( dwork_array, dwork, 1, 0, 0, dwork_msize, batchCount, queue );
magma_sset_pointer( dinvA_array, dinvA, STRTRI_BATCHED_NB, 0, 0, invA_msize, batchCount, queue );
magma_int_t streamid;
const magma_int_t nbstreams=10;
magma_queue_t queues[nbstreams];
for (k=0; k < nbstreams; k++) {
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[k] );
}
magma_getvector( batchCount, sizeof(float*), dA_array, 1, cpuAarray, 1, queue);
if (uplo == MagmaUpper) {
printf("Upper side is unavailable\n");
goto fin;
}
else {
for (j = 0; j < n; j += nb) {
ib = min(nb, n-j);
#if 1
//===============================================
// panel factorization
//===============================================
magma_sdisplace_pointers(dA_displ, dA_array, ldda, j, j, batchCount, queue);
magma_sset_pointer( dwork_array, dwork, 1, 0, 0, dwork_msize, batchCount, queue );
magma_sset_pointer( dinvA_array, dinvA, STRTRI_BATCHED_NB, 0, 0, invA_msize, batchCount, queue );
if (recnb == nb)
{
arginfo = magma_spotrf_panel_batched(
uplo, n-j, ib,
dA_displ, ldda,
dwork_array, dwork_msize,
dinvA_array, invA_msize,
dW0_displ, dW1_displ, dW2_displ,
dW3_displ, dW4_displ,
info_array, j, batchCount, queue);
}
else {
//arginfo = magma_spotrf_rectile_batched(
arginfo = magma_spotrf_recpanel_batched(
uplo, n-j, ib, recnb,
dA_displ, ldda,
dwork_array, dwork_msize,
dinvA_array, invA_msize,
dW0_displ, dW1_displ, dW2_displ,
dW3_displ, dW4_displ,
info_array, j, batchCount, queue);
}
if (arginfo != 0 ) goto fin;
//===============================================
// end of panel
//===============================================
#endif
#if 1
//real_Double_t gpu_time;
//gpu_time = magma_sync_wtime(queue);
if ( (n-j-ib) > 0) {
use_stream = magma_srecommend_cublas_gemm_stream(MagmaNoTrans, MagmaConjTrans, n-j-ib, n-j-ib, ib);
if (use_stream)
{
//-------------------------------------------
// USE STREAM HERK
//-------------------------------------------
// since it use different queue I need to wait the panel.
/* you must know the matrix layout inorder to do it */
magma_queue_sync(queue);
for (k=0; k < batchCount; k++)
{
streamid = k%nbstreams;
// call herk, class ssyrk must call cpu pointer
magma_ssyrk( MagmaLower, MagmaNoTrans, n-j-ib, ib,
d_alpha,
(const float*) cpuAarray[k] + j+ib+j*ldda, ldda,
d_beta,
cpuAarray[k] + j+ib+(j+ib)*ldda, ldda, queues[streamid] );
}
// need to synchronise to be sure that panel do not start before
// finishing the update at least of the next panel
// if queue is NULL, no need to sync
if ( queue != NULL ) {
for (magma_int_t s=0; s < nbstreams; s++)
magma_queue_sync(queues[s]);
}
}
else
{
//-------------------------------------------
// USE BATCHED GEMM(which is a HERK in fact, since it only access the lower part)
//-------------------------------------------
magma_sdisplace_pointers(dA_displ, dA_array, ldda, j+ib, j, batchCount, queue);
magma_sdisplace_pointers(dW1_displ, dA_array, ldda, j+ib, j+ib, batchCount, queue);
magmablas_ssyrk_batched( uplo, MagmaNoTrans, n-j-ib, ib,
d_alpha, dA_displ, ldda,
d_beta, dW1_displ, ldda,
batchCount, queue );
}
}
//gpu_time = magma_sync_wtime(queue) - gpu_time;
//real_Double_t flops = (n-j-ib) * (n-j-ib) * ib / 1e9 * batchCount;
//real_Double_t gpu_perf = flops / gpu_time;
//printf("Rows= %lld, Colum=%lld, herk time = %7.2fms, Gflops= %7.2f\n",
// (long long)(n-j-ib), (long long) ib, gpu_time*1000, gpu_perf);
#endif
}
}
fin:
magma_queue_sync(queue);
for (k=0; k < nbstreams; k++) {
magma_queue_destroy( queues[k] );
}
magma_free(dA_displ);
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dinvA_array);
magma_free(dwork_array);
magma_free( dinvA );
magma_free( dwork );
magma_free_cpu(cpuAarray);
return arginfo;
}
