/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgetri_batched
*/
int main( int argc, char** argv)
{
TESTING_INIT();
// constants
const magmaDoubleComplex c_zero = MAGMA_Z_ZERO;
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
const magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
magmaDoubleComplex *h_A, *h_Ainv, *h_R, *work;
magmaDoubleComplex_ptr d_A, d_invA;
magmaDoubleComplex_ptr *dA_array;
magmaDoubleComplex_ptr *dinvA_array;
magma_int_t **dipiv_array;
magma_int_t *dinfo_array;
magma_int_t *ipiv, *cpu_info;
magma_int_t *d_ipiv, *d_info;
magma_int_t N, n2, lda, ldda, info, info1, info2, lwork;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magmaDoubleComplex tmp;
double error, rwork[1];
magma_int_t columns;
magma_int_t status = 0;
magma_opts opts( MagmaOptsBatched );
opts.parse_opts( argc, argv );
magma_int_t batchCount = opts.batchcount;
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("%% batchCount N CPU Gflop/s (ms) GPU Gflop/s (ms) ||I - A*A^{-1}||_1 / (N*cond(A))\n");
printf("%%===============================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
N = opts.nsize[itest];
lda = N;
n2 = lda*N * batchCount;
ldda = magma_roundup( N, opts.align ); // multiple of 32 by default
// This is the correct flops but since this getri_batched is based on
// 2 trsm = getrs and to know the real flops I am using the getrs one
//gflops = (FLOPS_ZGETRF( N, N ) + FLOPS_ZGETRI( N ))/ 1e9 * batchCount;
gflops = (FLOPS_ZGETRF( N, N ) + FLOPS_ZGETRS( N, N ))/ 1e9 * batchCount;
// query for workspace size
lwork = -1;
lapackf77_zgetri( &N, NULL, &lda, NULL, &tmp, &lwork, &info );
if (info != 0) {
printf("lapackf77_zgetri returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
lwork = magma_int_t( MAGMA_Z_REAL( tmp ));
TESTING_MALLOC_CPU( cpu_info, magma_int_t, batchCount );
TESTING_MALLOC_CPU( ipiv, magma_int_t, N * batchCount );
TESTING_MALLOC_CPU( work, magmaDoubleComplex, lwork*batchCount );
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_CPU( h_Ainv, magmaDoubleComplex, n2 );
TESTING_MALLOC_CPU( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N * batchCount );
TESTING_MALLOC_DEV( d_invA, magmaDoubleComplex, ldda*N * batchCount );
TESTING_MALLOC_DEV( d_ipiv, magma_int_t, N * batchCount );
TESTING_MALLOC_DEV( d_info, magma_int_t, batchCount );
TESTING_MALLOC_DEV( dA_array, magmaDoubleComplex*, batchCount );
TESTING_MALLOC_DEV( dinvA_array, magmaDoubleComplex*, batchCount );
TESTING_MALLOC_DEV( dinfo_array, magma_int_t, batchCount );
TESTING_MALLOC_DEV( dipiv_array, magma_int_t*, batchCount );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
columns = N * batchCount;
lapackf77_zlacpy( MagmaFullStr, &N, &columns, h_A, &lda, h_R, &lda );
lapackf77_zlacpy( MagmaFullStr, &N, &columns, h_A, &lda, h_Ainv, &lda );
magma_zsetmatrix( N, columns, h_R, lda, d_A, ldda, opts.queue );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zset_pointer( dA_array, d_A, ldda, 0, 0, ldda * N, batchCount, opts.queue );
magma_zset_pointer( dinvA_array, d_invA, ldda, 0, 0, ldda * N, batchCount, opts.queue );
magma_iset_pointer( dipiv_array, d_ipiv, 1, 0, 0, N, batchCount, opts.queue );
gpu_time = magma_sync_wtime( opts.queue );
info1 = magma_zgetrf_batched( N, N, dA_array, ldda, dipiv_array, dinfo_array, batchCount, opts.queue);
info2 = magma_zgetri_outofplace_batched( N, dA_array, ldda, dipiv_array, dinvA_array, ldda, dinfo_array, batchCount, opts.queue);
gpu_time = magma_sync_wtime( opts.queue ) - gpu_time;
gpu_perf = gflops / gpu_time;
// check correctness of results throught "dinfo_magma" and correctness of argument throught "info"
magma_getvector( batchCount, sizeof(magma_int_t), dinfo_array, 1, cpu_info, 1, opts.queue );
for (magma_int_t i=0; i < batchCount; i++)
{
if (cpu_info[i] != 0 ) {
printf("magma_zgetrf_batched matrix %d returned error %d\n", (int) i, (int)cpu_info[i] );
}
}
if (info1 != 0) printf("magma_zgetrf_batched returned argument error %d: %s.\n", (int) info1, magma_strerror( info1 ));
if (info2 != 0) printf("magma_zgetri_batched returned argument error %d: %s.\n", (int) info2, magma_strerror( info2 ));
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
#if !defined (BATCHED_DISABLE_PARCPU) && defined(_OPENMP)
magma_int_t nthreads = magma_get_lapack_numthreads();
magma_set_lapack_numthreads(1);
magma_set_omp_numthreads(nthreads);
#pragma omp parallel for schedule(dynamic)
#endif
for (int i=0; i < batchCount; i++)
{
magma_int_t locinfo;
lapackf77_zgetrf(&N, &N, h_Ainv + i*lda*N, &lda, ipiv + i*N, &locinfo);
if (locinfo != 0) {
printf("lapackf77_zgetrf returned error %d: %s.\n",
(int) locinfo, magma_strerror( locinfo ));
}
lapackf77_zgetri(&N, h_Ainv + i*lda*N, &lda, ipiv + i*N, work + i*lwork, &lwork, &locinfo );
if (locinfo != 0) {
printf("lapackf77_zgetri returned error %d: %s.\n",
(int) locinfo, magma_strerror( locinfo ));
}
}
#if !defined (BATCHED_DISABLE_PARCPU) && defined(_OPENMP)
magma_set_lapack_numthreads(nthreads);
#endif
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
printf("%10d %5d %7.2f (%7.2f) %7.2f (%7.2f)",
(int) batchCount, (int) N, cpu_perf, cpu_time*1000., gpu_perf, gpu_time*1000. );
}
else {
printf("%10d %5d --- ( --- ) %7.2f (%7.2f)",
(int) batchCount, (int) N, gpu_perf, gpu_time*1000. );
}
/* =====================================================================
Check the result
=================================================================== */
if ( opts.check ) {
magma_igetvector( N*batchCount, d_ipiv, 1, ipiv, 1, opts.queue );
magma_zgetmatrix( N, N*batchCount, d_invA, ldda, h_Ainv, lda, opts.queue );
error = 0;
for (magma_int_t i=0; i < batchCount; i++)
{
for (magma_int_t k=0; k < N; k++) {
if (ipiv[i*N+k] < 1 || ipiv[i*N+k] > N )
{
printf("error for matrix %d ipiv @ %d = %d\n", (int) i, (int) k, (int) ipiv[i*N+k]);
error = -1;
}
}
if (error == -1) {
break;
}
// compute 1-norm condition number estimate, following LAPACK's zget03
double normA, normAinv, rcond, err;
normA = lapackf77_zlange( "1", &N, &N, h_A + i*lda*N, &lda, rwork );
normAinv = lapackf77_zlange( "1", &N, &N, h_Ainv + i*lda*N, &lda, rwork );
if ( normA <= 0 || normAinv <= 0 ) {
rcond = 0;
err = 1 / (tol/opts.tolerance); // == 1/eps
}
else {
rcond = (1 / normA) / normAinv;
// R = I
// R -= A*A^{-1}
// err = ||I - A*A^{-1}|| / ( N ||A||*||A^{-1}|| ) = ||R|| * rcond / N, using 1-norm
lapackf77_zlaset( "full", &N, &N, &c_zero, &c_one, h_R + i*lda*N, &lda );
blasf77_zgemm( "no", "no", &N, &N, &N, &c_neg_one,
h_A + i*lda*N, &lda,
h_Ainv + i*lda*N, &lda, &c_one,
h_R + i*lda*N, &lda );
err = lapackf77_zlange( "1", &N, &N, h_R + i*lda*N, &lda, rwork );
err = err * rcond / N;
}
if ( isnan(err) || isinf(err) ) {
error = err;
break;
}
error = max( err, error );
}
bool okay = (error < tol);
status += ! okay;
printf(" %8.2e %s\n", error, (okay ? "ok" : "failed") );
}
else {
printf("\n");
}
TESTING_FREE_CPU( cpu_info );
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( work );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_Ainv );
TESTING_FREE_CPU( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_invA );
TESTING_FREE_DEV( d_ipiv );
TESTING_FREE_DEV( d_info );
TESTING_FREE_DEV( dA_array );
TESTING_FREE_DEV( dinvA_array );
TESTING_FREE_DEV( dinfo_array );
TESTING_FREE_DEV( dipiv_array );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgetrf_batched
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cublas_perf=0., cublas_time=0., cpu_perf=0, cpu_time=0;
double error;
magma_int_t cublas_enable = 0;
magmaDoubleComplex *h_A, *h_R;
magmaDoubleComplex *dA_magma;
magmaDoubleComplex **dA_array = NULL;
magma_int_t **dipiv_array = NULL;
magma_int_t *ipiv, *cpu_info;
magma_int_t *dipiv_magma, *dinfo_magma;
magma_int_t M, N, n2, lda, ldda, min_mn, info;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t batchCount;
magma_int_t status = 0;
magma_opts opts( MagmaOptsBatched );
opts.parse_opts( argc, argv );
//opts.lapack |= opts.check;
batchCount = opts.batchcount;
magma_int_t columns;
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("%% BatchCount M N CPU Gflop/s (ms) MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) ||PA-LU||/(||A||*N)\n");
printf("%%==========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
n2 = lda*N * batchCount;
ldda = magma_roundup( M, opts.align ); // multiple of 32 by default
gflops = FLOPS_ZGETRF( M, N ) / 1e9 * batchCount;
TESTING_MALLOC_CPU( cpu_info, magma_int_t, batchCount );
TESTING_MALLOC_CPU( ipiv, magma_int_t, min_mn * batchCount );
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_CPU( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC_DEV( dA_magma, magmaDoubleComplex, ldda*N * batchCount );
TESTING_MALLOC_DEV( dipiv_magma, magma_int_t, min_mn * batchCount );
TESTING_MALLOC_DEV( dinfo_magma, magma_int_t, batchCount );
TESTING_MALLOC_DEV( dA_array, magmaDoubleComplex*, batchCount );
TESTING_MALLOC_DEV( dipiv_array, magma_int_t*, batchCount );
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
// make A diagonally dominant, to not need pivoting
for( int s=0; s < batchCount; ++s ) {
for( int i=0; i < min_mn; ++i ) {
h_A[ i + i*lda + s*lda*N ] = MAGMA_Z_MAKE(
MAGMA_Z_REAL( h_A[ i + i*lda + s*lda*N ] ) + N,
MAGMA_Z_IMAG( h_A[ i + i*lda + s*lda*N ] ));
}
}
columns = N * batchCount;
lapackf77_zlacpy( MagmaFullStr, &M, &columns, h_A, &lda, h_R, &lda );
magma_zsetmatrix( M, columns, h_R, lda, dA_magma, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zset_pointer( dA_array, dA_magma, ldda, 0, 0, ldda*N, batchCount, opts.queue );
magma_time = magma_sync_wtime( opts.queue );
info = magma_zgetrf_nopiv_batched( M, N, dA_array, ldda, dinfo_magma, batchCount, opts.queue);
magma_time = magma_sync_wtime( opts.queue ) - magma_time;
magma_perf = gflops / magma_time;
// check correctness of results throught "dinfo_magma" and correctness of argument throught "info"
magma_getvector( batchCount, sizeof(magma_int_t), dinfo_magma, 1, cpu_info, 1);
for (int i=0; i < batchCount; i++)
{
if (cpu_info[i] != 0 ) {
printf("magma_zgetrf_batched matrix %d returned internal error %d\n", i, (int)cpu_info[i] );
}
}
if (info != 0) {
printf("magma_zgetrf_batched returned argument error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
for (int i=0; i < batchCount; i++) {
lapackf77_zgetrf(&M, &N, h_A + i*lda*N, &lda, ipiv + i * min_mn, &info);
assert( info == 0 );
}
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_zgetrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
}
/* =====================================================================
Check the factorization
=================================================================== */
if ( opts.lapack ) {
printf("%10d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f)",
(int) batchCount, (int) M, (int) N, cpu_perf, cpu_time*1000., magma_perf, magma_time*1000., cublas_perf*cublas_enable, cublas_time*1000.*cublas_enable );
}
else {
printf("%10d %5d %5d --- ( --- ) %7.2f (%7.2f) %7.2f (%7.2f)",
(int) batchCount, (int) M, (int) N, magma_perf, magma_time*1000., cublas_perf*cublas_enable, cublas_time*1000.*cublas_enable );
}
if ( opts.check ) {
// initialize ipiv to 1, 2, 3, ...
for (int i=0; i < batchCount; i++)
{
for (int k=0; k < min_mn; k++) {
ipiv[i*min_mn+k] = k+1;
}
}
magma_zgetmatrix( M, N*batchCount, dA_magma, ldda, h_A, lda );
error = 0;
for (int i=0; i < batchCount; i++)
{
double err;
err = get_LU_error( M, N, h_R + i * lda*N, lda, h_A + i * lda*N, ipiv + i * min_mn);
if ( isnan(err) || isinf(err) ) {
error = err;
break;
}
error = max( err, error );
}
bool okay = (error < tol);
status += ! okay;
printf(" %8.2e %s\n", error, (okay ? "ok" : "failed") );
}
else {
printf(" --- \n");
}
TESTING_FREE_CPU( cpu_info );
TESTING_FREE_CPU( ipiv );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_R );
TESTING_FREE_DEV( dA_magma );
TESTING_FREE_DEV( dinfo_magma );
TESTING_FREE_DEV( dipiv_magma );
TESTING_FREE_DEV( dipiv_array );
TESTING_FREE_DEV( dA_array );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
extern "C" magma_int_t
magma_zgeqrf_panel_batched(
magma_int_t m, magma_int_t n, magma_int_t nb,
magmaDoubleComplex** dA_array, magma_int_t ldda,
magmaDoubleComplex** tau_array,
magmaDoubleComplex** dT_array, magma_int_t ldt,
magmaDoubleComplex** dR_array, magma_int_t ldr,
magmaDoubleComplex** dW0_displ,
magmaDoubleComplex** dW1_displ,
magmaDoubleComplex *dwork,
magmaDoubleComplex** dW2_displ,
magmaDoubleComplex** dW3_displ,
magma_int_t *info_array,
magma_int_t batchCount, magma_queue_t queue)
{
magma_int_t j, jb;
magma_int_t ldw = nb;
magma_int_t minmn = min(m,n);
for( j=0; j < minmn; j += nb)
{
jb = min(nb, minmn-j);
magma_zdisplace_pointers(dW0_displ, dA_array, ldda, j, j, batchCount, queue);
magma_zdisplace_pointers(dW2_displ, tau_array, 1, j, 0, batchCount, queue);
magma_zdisplace_pointers(dW3_displ, dR_array, ldr, j, j, batchCount, queue); //
//sub-panel factorization
magma_zgeqr2_batched(
m-j, jb,
dW0_displ, ldda,
dW2_displ,
info_array,
batchCount,
queue);
//copy th whole rectangular n,jb from of dA to dR (it's lower portion (which is V's) will be set to zero if needed at the end)
magma_zdisplace_pointers(dW0_displ, dA_array, ldda, 0, j, batchCount, queue);
magma_zdisplace_pointers(dW3_displ, dR_array, ldr, 0, j, batchCount, queue);
magmablas_zlacpy_batched( MagmaFull, minmn, jb, dW0_displ, ldda, dW3_displ, ldr, batchCount, queue );
//set the upper jbxjb portion of V dA(j,j) to 1/0s (note that the rectangular on the top of this triangular of V still non zero but has been copied to dR).
magma_zdisplace_pointers(dW0_displ, dA_array, ldda, j, j, batchCount, queue);
magmablas_zlaset_batched( MagmaUpper, jb, jb, MAGMA_Z_ZERO, MAGMA_Z_ONE, dW0_displ, ldda, batchCount, queue );
if ( (n-j-jb) > 0) //update the trailing matrix inside the panel
{
magma_zlarft_sm32x32_batched(m-j, jb,
dW0_displ, ldda,
dW2_displ,
dT_array, ldt,
batchCount, queue);
magma_zdisplace_pointers( dW1_displ, dA_array, ldda, j, j + jb, batchCount, queue );
magma_zset_pointer( dW2_displ, dwork, 1, 0, 0, ldw*n, batchCount, queue );
magma_zset_pointer( dW3_displ, dwork + ldw*n*batchCount, 1, 0, 0, ldw*n, batchCount, queue );
magma_zlarfb_gemm_batched(
MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-j, n-j-jb, jb,
(const magmaDoubleComplex**)dW0_displ, ldda,
(const magmaDoubleComplex**)dT_array, ldt,
dW1_displ, ldda,
dW2_displ, ldw,
dW3_displ, ldw,
batchCount, queue );
}
}
// copy the remaining portion of dR from dA in case m < n
if ( m < n )
{
magma_zdisplace_pointers(dW0_displ, dA_array, ldda, 0, minmn, batchCount, queue);
magma_zdisplace_pointers(dW3_displ, dR_array, ldr, 0, minmn, batchCount, queue);
magmablas_zlacpy_batched( MagmaFull, minmn, n-minmn, dW0_displ, ldda, dW3_displ, ldr, batchCount, queue );
}
// to be consistent set the whole upper nbxnb of V to 0/1s, in this case no need to set it inside zgeqrf_batched
magma_zdisplace_pointers(dW0_displ, dA_array, ldda, 0, 0, batchCount, queue);
magmablas_zlaset_batched( MagmaUpper, minmn, n, MAGMA_Z_ZERO, MAGMA_Z_ONE, dW0_displ, ldda, batchCount, queue );
return MAGMA_SUCCESS;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing zgeqrf_batched
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, magma_perf, magma_time, cublas_perf=0, cublas_time=0, cpu_perf, cpu_time;
double magma_error, cublas_error, magma_error2, cublas_error2;
magmaDoubleComplex *h_A, *h_R, *h_Amagma, *tau, *h_work, tmp[1];
magmaDoubleComplex *d_A, *dtau_magma, *dtau_cublas;
magmaDoubleComplex **dA_array = NULL;
magmaDoubleComplex **dtau_array = NULL;
magma_int_t *dinfo_magma, *dinfo_cublas;
magma_int_t M, N, lda, ldda, lwork, n2, info, min_mn;
magma_int_t ione = 1;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t status = 0;
magma_int_t batchCount;
magma_int_t column;
magma_opts opts( MagmaOptsBatched );
opts.parse_opts( argc, argv );
batchCount = opts.batchcount;
double tol = opts.tolerance * lapackf77_dlamch("E");
printf("%% BatchCount M N MAGMA Gflop/s (ms) CUBLAS Gflop/s (ms) CPU Gflop/s (ms) |R - Q^H*A|_mag |I - Q^H*Q|_mag |R - Q^H*A|_cub |I - Q^H*Q|_cub\n");
printf("%%============================================================================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
min_mn = min(M, N);
lda = M;
n2 = lda*N * batchCount;
ldda = M;
ldda = magma_roundup( M, opts.align ); // multiple of 32 by default
gflops = (FLOPS_ZGEQRF( M, N ) + FLOPS_ZGEQRT( M, N )) / 1e9 * batchCount;
/* Allocate memory for the matrix */
TESTING_MALLOC_CPU( tau, magmaDoubleComplex, min_mn * batchCount );
TESTING_MALLOC_CPU( h_A, magmaDoubleComplex, n2 );
TESTING_MALLOC_CPU( h_Amagma, magmaDoubleComplex, n2 );
TESTING_MALLOC_PIN( h_R, magmaDoubleComplex, n2 );
TESTING_MALLOC_DEV( d_A, magmaDoubleComplex, ldda*N * batchCount );
TESTING_MALLOC_DEV( dtau_magma, magmaDoubleComplex, min_mn * batchCount);
TESTING_MALLOC_DEV( dtau_cublas, magmaDoubleComplex, min_mn * batchCount);
TESTING_MALLOC_DEV( dinfo_magma, magma_int_t, batchCount);
TESTING_MALLOC_DEV( dinfo_cublas, magma_int_t, batchCount);
TESTING_MALLOC_DEV( dA_array, magmaDoubleComplex*, batchCount );
TESTING_MALLOC_DEV( dtau_array, magmaDoubleComplex*, batchCount );
// to determine the size of lwork
lwork = -1;
lapackf77_zgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_Z_REAL( tmp[0] );
lwork = max(lwork, N*N);
TESTING_MALLOC_CPU( h_work, magmaDoubleComplex, lwork * batchCount);
column = N * batchCount;
/* Initialize the matrix */
lapackf77_zlarnv( &ione, ISEED, &n2, h_A );
lapackf77_zlacpy( MagmaFullStr, &M, &column, h_A, &lda, h_R, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zsetmatrix( M, column, h_R, lda, d_A, ldda );
magma_zset_pointer( dA_array, d_A, 1, 0, 0, ldda*N, batchCount, opts.queue );
magma_zset_pointer( dtau_array, dtau_magma, 1, 0, 0, min_mn, batchCount, opts.queue );
magma_time = magma_sync_wtime( opts.queue );
info = magma_zgeqrf_batched(M, N, dA_array, ldda, dtau_array, dinfo_magma, batchCount, opts.queue);
magma_time = magma_sync_wtime( opts.queue ) - magma_time;
magma_perf = gflops / magma_time;
magma_zgetmatrix( M, column, d_A, ldda, h_Amagma, lda);
if (info != 0) {
printf("magma_zgeqrf_batched returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* ====================================================================
Performs operation using CUBLAS
=================================================================== */
/* cublasZgeqrfBatched is only available from CUBLAS v6.5 */
#if CUDA_VERSION >= 6050
magma_zsetmatrix( M, column, h_R, lda, d_A, ldda );
magma_zset_pointer( dA_array, d_A, 1, 0, 0, ldda*N, batchCount, opts.queue );
magma_zset_pointer( dtau_array, dtau_cublas, 1, 0, 0, min_mn, batchCount, opts.queue );
cublas_time = magma_sync_wtime( opts.queue );
int cublas_info; // not magma_int_t
cublasZgeqrfBatched( opts.handle, M, N, dA_array, ldda, dtau_array, &cublas_info, batchCount);
cublas_time = magma_sync_wtime( opts.queue ) - cublas_time;
cublas_perf = gflops / cublas_time;
if (cublas_info != 0) {
printf("cublasZgeqrfBatched returned error %d: %s.\n",
(int) cublas_info, magma_strerror( cublas_info ));
}
#endif
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.check ) {
cpu_time = magma_wtime();
// #define BATCHED_DISABLE_PARCPU
#if !defined (BATCHED_DISABLE_PARCPU) && defined(_OPENMP)
magma_int_t nthreads = magma_get_lapack_numthreads();
magma_set_lapack_numthreads(1);
magma_set_omp_numthreads(nthreads);
#pragma omp parallel for schedule(dynamic)
#endif
for (magma_int_t s=0; s < batchCount; s++)
{
magma_int_t locinfo;
lapackf77_zgeqrf(&M, &N, h_A + s * lda * N, &lda, tau + s * min_mn, h_work + s * lwork, &lwork, &locinfo);
if (locinfo != 0) {
printf("lapackf77_zgeqrf matrix %d returned error %d: %s.\n",
(int) s, (int) locinfo, magma_strerror( locinfo ));
}
}
#if !defined (BATCHED_DISABLE_PARCPU) && defined(_OPENMP)
magma_set_lapack_numthreads(nthreads);
#endif
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0) {
printf("lapackf77_zgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Check the MAGMA CUBLAS result compared to LAPACK
=================================================================== */
magma_int_t ldq = M;
magma_int_t ldr = min_mn;
magmaDoubleComplex *Q, *R;
double *work;
TESTING_MALLOC_CPU( Q, magmaDoubleComplex, ldq*min_mn ); // M by K
TESTING_MALLOC_CPU( R, magmaDoubleComplex, ldr*N ); // K by N
TESTING_MALLOC_CPU( work, double, min_mn );
/* check magma result */
magma_error = 0;
magma_error2 = 0;
magma_zgetvector(min_mn*batchCount, dtau_magma, 1, tau, 1);
for (int i=0; i < batchCount; i++)
{
double err, err2;
get_QR_error(M, N, min_mn,
h_Amagma + i*lda*N, h_R + i*lda*N, lda, tau + i*min_mn,
Q, ldq, R, ldr, h_work, lwork,
work, &err, &err2);
if ( isnan(err) || isinf(err) ) {
magma_error = err;
break;
}
magma_error = max( err, magma_error );
magma_error2 = max( err2, magma_error2 );
}
/* check cublas result */
cublas_error = 0;
cublas_error2 = 0;
#if CUDA_VERSION >= 6050
magma_zgetvector(min_mn*batchCount, dtau_magma, 1, tau, 1);
magma_zgetmatrix( M, column, d_A, ldda, h_A, lda);
for (int i=0; i < batchCount; i++)
{
double err, err2;
get_QR_error(M, N, min_mn,
h_A + i*lda*N, h_R + i*lda*N, lda, tau + i*min_mn,
Q, ldq, R, ldr, h_work, lwork,
work, &err, &err2);
if ( isnan(err) || isinf(err) ) {
cublas_error = err;
break;
}
cublas_error = max( err, cublas_error );
cublas_error2 = max( err2, cublas_error2 );
}
#endif
TESTING_FREE_CPU( Q ); Q = NULL;
TESTING_FREE_CPU( R ); R = NULL;
TESTING_FREE_CPU( work ); work = NULL;
bool okay = (magma_error < tol && magma_error2 < tol);
//bool okay_cublas = (cublas_error < tol && cublas_error2 < tol);
status += ! okay;
printf("%10d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %7.2f (%7.2f) %15.2e %15.2e %15.2e %15.2e %s\n",
(int)batchCount, (int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time,
cpu_perf, 1000.*cpu_time,
magma_error, magma_error2,
cublas_error, cublas_error2,
(okay ? "ok" : "failed") );
}
else {
printf("%10d %5d %5d %7.2f (%7.2f) %7.2f (%7.2f) --- ( --- ) ---\n",
(int)batchCount, (int) M, (int) N,
magma_perf, 1000.*magma_time,
cublas_perf, 1000.*cublas_time );
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_Amagma);
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( dtau_magma );
TESTING_FREE_DEV( dtau_cublas );
TESTING_FREE_DEV( dinfo_magma );
TESTING_FREE_DEV( dinfo_cublas );
TESTING_FREE_DEV( dA_array );
TESTING_FREE_DEV( dtau_array );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
