extern "C" magma_int_t
magma_dorgqr(
magma_int_t m, magma_int_t n, magma_int_t k,
double *a, magma_int_t lda,
double *tau, magmaDouble_ptr dT, size_t dT_offset,
magma_int_t nb,
magma_queue_t queue,
magma_int_t *info )
{
/* -- clMAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
@date November 2014
Purpose
=======
DORGQR generates an M-by-N DOUBLE_PRECISION matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by DGEQRF.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix Q. M >= 0.
N (input) INTEGER
The number of columns of the matrix Q. M >= N >= 0.
K (input) INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
A (input/output) DOUBLE_PRECISION array A, dimension (LDDA,N).
On entry, the i-th column must contain the vector
which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by DGEQRF_GPU in the
first k columns of its array argument A.
On exit, the M-by-N matrix Q.
LDA (input) INTEGER
The first dimension of the array A. LDA >= max(1,M).
TAU (input) DOUBLE_PRECISION array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by DGEQRF_GPU.
DT (input) DOUBLE_PRECISION array on the GPU device.
DT contains the T matrices used in blocking the elementary
reflectors H(i), e.g., this can be the 6th argument of
magma_dgeqrf_gpu.
NB (input) INTEGER
This is the block size used in DGEQRF_GPU, and correspondingly
the size of the T matrices, used in the factorization, and
stored in DT.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument has an illegal value
===================================================================== */
#define a_ref(i,j) ( a + (j)*lda + (i))
#define da_ref(i,j) da, (da_offset + (j)*ldda + (i))
#define t_ref(a_1) dT, (dT_offset + (a_1)*nb)
double c_zero = MAGMA_D_ZERO;
magma_int_t i__1, i__2, i__3;
magma_int_t lwork, ldda;
magma_int_t i, ib, ki, kk, iinfo;
magma_int_t lddwork = min(m, n);
double *work;
magmaDouble_ptr da, dwork;
size_t da_offset, dwork_offset;
magma_event_t event = NULL;
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (lda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (n <= 0)
return *info;
/* Allocate GPU work space */
ldda = ((m+31)/32)*32;
lddwork = ((lddwork+31)/32)*32;
if (MAGMA_SUCCESS != magma_dmalloc( &da, ((n)*ldda + nb*lddwork ) )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
da_offset = 0;
dwork = da;
dwork_offset = da_offset + (n)*ldda;
/* Allocate CPU work space */
lwork = n * nb;
magma_dmalloc_cpu( &work, lwork );
if( work == NULL ) {
magma_free( da );
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
if ( (nb > 1) && (nb < k) )
{
/* Use blocked code after the last block.
The first kk columns are handled by the block method. */
ki = (k - nb - 1) / nb * nb;
kk = min(k, ki + nb);
/* Set A(1:kk,kk+1:n) to zero. */
magmablas_dlaset(MagmaFull, kk, n-kk, c_zero, c_zero, da_ref(0,kk), ldda, queue);
}
else
kk = 0;
/* Use unblocked code for the last or only block. */
if (kk < n)
{
i__1 = m - kk;
i__2 = n - kk;
i__3 = k - kk;
lapackf77_dorgqr(&i__1, &i__2, &i__3,
a_ref(kk, kk), &lda,
&tau[kk], work, &lwork, &iinfo);
magma_dsetmatrix(i__1, i__2, a_ref(kk, kk), lda, da_ref(kk, kk), ldda, queue);
}
if (kk > 0)
{
/* Use blocked code */
for (i = ki; i >= 0; i-=nb)
{
ib = min(nb, k - i);
/* Send the current panel to the GPU */
i__2 = m - i;
dpanel_to_q(MagmaUpper, ib, a_ref(i,i), lda, work);
magma_dsetmatrix(i__2, ib, a_ref(i, i), lda, da_ref(i, i), ldda, queue);
if (i + ib < n)
{
/* Apply H to A(i:m,i+ib:n) from the left */
i__3 = n - i - ib;
magma_dlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
i__2, i__3, ib,
da_ref(i, i ), ldda, t_ref(i), nb,
da_ref(i, i+ib), ldda, dwork, dwork_offset, lddwork, queue);
}
/* Apply H to rows i:m of current block on the CPU */
lapackf77_dorgqr(&i__2, &ib, &ib,
a_ref(i, i), &lda,
&tau[i], work, &lwork, &iinfo);
magma_dsetmatrix_async( i__2, ib,
a_ref(i,i), lda,
da_ref(i,i), ldda, queue, &event );
/* Set rows 1:i-1 of current block to zero */
i__2 = i + ib;
magmablas_dlaset(MagmaFull, i, i__2 - i, c_zero, c_zero, da_ref(0,i), ldda, queue);
}
}
magma_dgetmatrix(m, n, da_ref(0, 0), ldda, a_ref(0, 0), lda, queue);
//cudaStreamDestroy(stream);
magma_free( da );
magma_free_cpu(work);
return *info;
} /* magma_dorgqr */
extern "C" magma_int_t
magma_dorgqr_gpu(magma_int_t m, magma_int_t n, magma_int_t k,
double *dA, magma_int_t ldda,
double *tau,
double *dT, magma_int_t nb,
magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
DORGQR generates an M-by-N DOUBLE_PRECISION matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by DGEQRF_GPU.
Arguments
=========
M (input) INTEGER
The number of rows of the matrix Q. M >= 0.
N (input) INTEGER
The number of columns of the matrix Q. M >= N >= 0.
K (input) INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
DA (input/output) DOUBLE_PRECISION array A on the GPU, dimension (LDDA,N).
On entry, the i-th column must contain the vector
which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by DGEQRF_GPU in the
first k columns of its array argument A.
On exit, the M-by-N matrix Q.
LDDA (input) INTEGER
The first dimension of the array A. LDDA >= max(1,M).
TAU (input) DOUBLE_PRECISION array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by DGEQRF_GPU.
DT (input/workspace) DOUBLE_PRECISION work space array on the GPU,
dimension (2*MIN(M, N) + (N+31)/32*32 )*NB.
This must be the 6th argument of magma_dgeqrf_gpu
[ note that if N here is bigger than N in magma_dgeqrf_gpu,
the workspace requirement DT in magma_dgeqrf_gpu must be
as specified in this routine ].
NB (input) INTEGER
This is the block size used in DGEQRF_GPU, and correspondingly
the size of the T matrices, used in the factorization, and
stored in DT.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument has an illegal value
===================================================================== */
#define dA(i,j) (dA + (i) + (j)*ldda)
#define dT(j) (dT + (j)*nb)
magma_int_t m_kk, n_kk, k_kk, mi;
magma_int_t lwork, lpanel;
magma_int_t i, ib, ki, kk, iinfo;
magma_int_t lddwork;
double *dV, *dW;
double *work, *panel;
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (ldda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (n <= 0) {
return *info;
}
// first kk columns are handled by blocked method.
// ki is start of 2nd-to-last block
if ((nb > 1) && (nb < k)) {
ki = (k - nb - 1) / nb * nb;
kk = min( k, ki+nb );
} else {
ki = 0;
kk = 0;
}
// Allocate CPU work space
// n*nb for dorgqr workspace
// (m - kk)*(n - kk) for last block's panel
lwork = n*nb;
lpanel = (m - kk)*(n - kk);
magma_dmalloc_cpu( &work, lwork + lpanel );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
panel = work + lwork;
// Allocate work space on GPU
if (MAGMA_SUCCESS != magma_dmalloc( &dV, ldda*nb )) {
magma_free_cpu( work );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
// dT workspace has:
// 2*min(m,n)*nb for T and R^{-1} matrices from geqrf
// ((n+31)/32*32 )*nb for dW larfb workspace.
lddwork = min(m,n);
dW = dT + 2*lddwork*nb;
magma_queue_t stream;
magma_queue_create( &stream );
// Use unblocked code for the last or only block.
if (kk < n) {
m_kk = m - kk;
n_kk = n - kk;
k_kk = k - kk;
magma_dgetmatrix( m_kk, k_kk,
dA(kk, kk), ldda, panel, m_kk );
lapackf77_dorgqr( &m_kk, &n_kk, &k_kk,
panel, &m_kk,
&tau[kk], work, &lwork, &iinfo );
magma_dsetmatrix( m_kk, n_kk,
panel, m_kk, dA(kk, kk), ldda );
// Set A(1:kk,kk+1:n) to zero.
magmablas_dlaset( MagmaUpperLower, kk, n - kk, dA(0, kk), ldda );
}
if (kk > 0) {
// Use blocked code
// stream: copy Aii to V --> laset --> laset --> larfb --> [next]
// CPU has no computation
magmablasSetKernelStream( stream );
for (i = ki; i >= 0; i -= nb) {
ib = min( nb, k-i );
mi = m - i;
// Copy current panel on the GPU from dA to dV
magma_dcopymatrix_async( mi, ib,
dA(i,i), ldda,
dV, ldda, stream );
// set panel to identity
magmablas_dlaset( MagmaUpperLower, i, ib, dA(0, i), ldda );
magmablas_dlaset_identity( mi, ib, dA(i, i), ldda );
if (i < n) {
// Apply H to A(i:m,i:n) from the left
magma_dlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
mi, n-i, ib,
dV, ldda, dT(i), nb,
dA(i, i), ldda, dW, lddwork );
}
}
}
magma_queue_sync( stream );
magmablasSetKernelStream( NULL );
magma_free( dV );
magma_free_cpu( work );
magma_queue_destroy( stream );
return *info;
} /* magma_dorgqr_gpu */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgeqrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
const double d_neg_one = MAGMA_D_NEG_ONE;
const double d_one = MAGMA_D_ONE;
const double c_neg_one = MAGMA_D_NEG_ONE;
const double c_one = MAGMA_D_ONE;
const double c_zero = MAGMA_D_ZERO;
const magma_int_t ione = 1;
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf=0, cpu_time=0;
double Anorm, error=0, error2=0;
double *h_A, *h_R, *tau, *h_work, tmp[1];
magmaDouble_ptr d_A, dT;
magma_int_t M, N, n2, lda, ldda, lwork, info, min_mn, nb, size;
magma_int_t ISEED[4] = {0,0,0,1};
magma_opts opts;
parse_opts( argc, argv, &opts );
magma_int_t status = 0;
double tol = opts.tolerance * lapackf77_dlamch("E");
printf( "version %d\n", (int) opts.version );
if ( opts.version == 2 ) {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) |R - Q^H*A| |I - Q^H*Q|\n");
printf("===============================================================================\n");
}
else {
printf(" M N CPU GFlop/s (sec) GPU GFlop/s (sec) |b - A*x|\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;
ldda = ((M+31)/32)*32;
gflops = FLOPS_DGEQRF( M, N ) / 1e9;
// query for workspace size
lwork = -1;
lapackf77_dgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_D_REAL( tmp[0] );
TESTING_MALLOC_CPU( tau, double, min_mn );
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_CPU( h_work, double, lwork );
TESTING_MALLOC_PIN( h_R, double, n2 );
TESTING_MALLOC_DEV( d_A, double, ldda*N );
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_dsetmatrix( M, N, h_R, lda, d_A, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_wtime();
if ( opts.version == 2 ) {
// LAPACK complaint arguments
magma_dgeqrf2_gpu( M, N, d_A, ldda, tau, &info );
}
else {
nb = magma_get_dgeqrf_nb( M );
size = (2*min(M, N) + (N+31)/32*32 )*nb;
TESTING_MALLOC_DEV( dT, double, size );
if ( opts.version == 1 ) {
// stores dT, V blocks have zeros, R blocks inverted & stored in dT
magma_dgeqrf_gpu( M, N, d_A, ldda, tau, dT, &info );
}
#ifdef HAVE_CUBLAS
else if ( opts.version == 3 ) {
// stores dT, V blocks have zeros, R blocks stored in dT
magma_dgeqrf3_gpu( M, N, d_A, ldda, tau, dT, &info );
}
#endif
else {
printf( "Unknown version %d\n", (int) opts.version );
exit(1);
}
}
gpu_time = magma_wtime() - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0)
printf("magma_dgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
if ( opts.check && opts.version == 2 ) {
/* =====================================================================
Check the result, following zqrt01 except using the reduced Q.
This works for any M,N (square, tall, wide).
Only for version 2, which has LAPACK complaint output.
=================================================================== */
magma_dgetmatrix( M, N, d_A, ldda, h_R, lda );
magma_int_t ldq = M;
magma_int_t ldr = min_mn;
double *Q, *R;
double *work;
TESTING_MALLOC_CPU( Q, double, ldq*min_mn ); // M by K
TESTING_MALLOC_CPU( R, double, ldr*N ); // K by N
TESTING_MALLOC_CPU( work, double, min_mn );
// generate M by K matrix Q, where K = min(M,N)
lapackf77_dlacpy( "Lower", &M, &min_mn, h_R, &lda, Q, &ldq );
lapackf77_dorgqr( &M, &min_mn, &min_mn, Q, &ldq, tau, h_work, &lwork, &info );
assert( info == 0 );
// copy K by N matrix R
lapackf77_dlaset( "Lower", &min_mn, &N, &c_zero, &c_zero, R, &ldr );
lapackf77_dlacpy( "Upper", &min_mn, &N, h_R, &lda, R, &ldr );
// error = || R - Q^H*A || / (N * ||A||)
blasf77_dgemm( "Conj", "NoTrans", &min_mn, &N, &M,
&c_neg_one, Q, &ldq, h_A, &lda, &c_one, R, &ldr );
Anorm = lapackf77_dlange( "1", &M, &N, h_A, &lda, work );
error = lapackf77_dlange( "1", &min_mn, &N, R, &ldr, work );
if ( N > 0 && Anorm > 0 )
error /= (N*Anorm);
// set R = I (K by K identity), then R = I - Q^H*Q
// error = || I - Q^H*Q || / N
lapackf77_dlaset( "Upper", &min_mn, &min_mn, &c_zero, &c_one, R, &ldr );
blasf77_dsyrk( "Upper", "Conj", &min_mn, &M, &d_neg_one, Q, &ldq, &d_one, R, &ldr );
error2 = lapackf77_dlansy( "1", "Upper", &min_mn, R, &ldr, work );
if ( N > 0 )
error2 /= N;
TESTING_FREE_CPU( Q ); Q = NULL;
TESTING_FREE_CPU( R ); R = NULL;
TESTING_FREE_CPU( work ); work = NULL;
}
else if ( opts.check && M >= N ) {
/* =====================================================================
Check the result by solving consistent linear system, A*x = b.
Only for versions 1 & 3 with M >= N.
=================================================================== */
magma_int_t lwork;
double *x, *b, *hwork;
magmaDouble_ptr d_B;
const double c_zero = MAGMA_D_ZERO;
const double c_one = MAGMA_D_ONE;
const double c_neg_one = MAGMA_D_NEG_ONE;
const magma_int_t ione = 1;
// initialize RHS, b = A*random
TESTING_MALLOC_CPU( x, double, N );
TESTING_MALLOC_CPU( b, double, M );
lapackf77_dlarnv( &ione, ISEED, &N, x );
blasf77_dgemv( "Notrans", &M, &N, &c_one, h_A, &lda, x, &ione, &c_zero, b, &ione );
// copy to GPU
TESTING_MALLOC_DEV( d_B, double, M );
magma_dsetvector( M, b, 1, d_B, 1 );
if ( opts.version == 1 ) {
// allocate hwork
magma_dgeqrs_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, tmp, -1, &info );
lwork = (magma_int_t)MAGMA_D_REAL( tmp[0] );
TESTING_MALLOC_CPU( hwork, double, lwork );
// solve linear system
magma_dgeqrs_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, hwork, lwork, &info );
if (info != 0)
printf("magma_dgeqrs returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE_CPU( hwork );
}
#ifdef HAVE_CUBLAS
else if ( opts.version == 3 ) {
// allocate hwork
magma_dgeqrs3_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, tmp, -1, &info );
lwork = (magma_int_t)MAGMA_D_REAL( tmp[0] );
TESTING_MALLOC_CPU( hwork, double, lwork );
// solve linear system
magma_dgeqrs3_gpu( M, N, 1,
d_A, ldda, tau, dT,
d_B, M, hwork, lwork, &info );
if (info != 0)
printf("magma_dgeqrs3 returned error %d: %s.\n",
(int) info, magma_strerror( info ));
TESTING_FREE_CPU( hwork );
}
#endif
else {
printf( "Unknown version %d\n", (int) opts.version );
exit(1);
}
magma_dgetvector( N, d_B, 1, x, 1 );
// compute r = Ax - b, saved in b
blasf77_dgemv( "Notrans", &M, &N, &c_one, h_A, &lda, x, &ione, &c_neg_one, b, &ione );
// compute residual |Ax - b| / (n*|A|*|x|)
double norm_x, norm_A, norm_r, work[1];
norm_A = lapackf77_dlange( "F", &M, &N, h_A, &lda, work );
norm_r = lapackf77_dlange( "F", &M, &ione, b, &M, work );
norm_x = lapackf77_dlange( "F", &N, &ione, x, &N, work );
TESTING_FREE_CPU( x );
TESTING_FREE_CPU( b );
TESTING_FREE_DEV( d_B );
error = norm_r / (N * norm_A * norm_x);
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.lapack ) {
cpu_time = magma_wtime();
lapackf77_dgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_dgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
}
/* =====================================================================
Print performance and error.
=================================================================== */
printf("%5d %5d ", (int) M, (int) N );
if ( opts.lapack ) {
printf( "%7.2f (%7.2f)", cpu_perf, cpu_time );
}
else {
printf(" --- ( --- )" );
}
printf( " %7.2f (%7.2f) ", gpu_perf, gpu_time );
if ( opts.check ) {
if ( opts.version == 2 ) {
bool okay = (error < tol && error2 < tol);
status += ! okay;
printf( "%11.2e %11.2e %s\n", error, error2, (okay ? "ok" : "failed") );
}
else if ( M >= N ) {
bool okay = (error < tol);
status += ! okay;
printf( "%10.2e %s\n", error, (okay ? "ok" : "failed") );
}
else {
printf( "(error check only for M >= N)\n" );
}
}
else {
printf( " ---\n" );
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
if ( opts.version != 2 )
TESTING_FREE_DEV( dT );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
TESTING_FINALIZE();
return status;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing dgeqrf
*/
int main( int argc, char** argv)
{
TESTING_INIT();
real_Double_t gflops, gpu_perf, gpu_time, cpu_perf, cpu_time;
double error, error2;
double c_zero = MAGMA_D_ZERO;
double c_neg_one = MAGMA_D_NEG_ONE;
double c_one = MAGMA_D_ONE;
double *h_A, *h_T, *h_R, *tau, *h_work, tmp[1];
magmaDouble_ptr d_A, d_T, ddA, dtau;
magmaDouble_ptr d_A2, d_T2, ddA2, dtau2;
magmaDouble_ptr dwork, dwork2;
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;
#define BLOCK_SIZE 64
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = 10. * opts.tolerance * lapackf77_dlamch("E");
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
printf("version %d\n", (int) opts.version );
printf(" M N CPU GFlop/s (ms) GPU GFlop/s (ms) ||R - Q^H*A|| ||R_T||\n");
printf("=============================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
M = opts.msize[itest];
N = opts.nsize[itest];
if (N > 128) {
printf("%5d %5d skipping because dgeqr2x requires N <= 128\n",
(int) M, (int) N);
continue;
}
if (M < N) {
printf("%5d %5d skipping because dgeqr2x requires M >= N\n",
(int) M, (int) N);
continue;
}
min_mn = min(M, N);
lda = M;
n2 = lda*N;
ldda = ((M+31)/32)*32;
gflops = (FLOPS_DGEQRF( M, N ) + FLOPS_DGEQRT( M, N )) / 1e9;
/* Allocate memory for the matrix */
TESTING_MALLOC_CPU( tau, double, min_mn );
TESTING_MALLOC_CPU( h_A, double, n2 );
TESTING_MALLOC_CPU( h_T, double, N*N );
TESTING_MALLOC_PIN( h_R, double, n2 );
TESTING_MALLOC_DEV( d_A, double, ldda*N );
TESTING_MALLOC_DEV( d_T, double, N*N );
TESTING_MALLOC_DEV( ddA, double, N*N );
TESTING_MALLOC_DEV( dtau, double, min_mn );
TESTING_MALLOC_DEV( d_A2, double, ldda*N );
TESTING_MALLOC_DEV( d_T2, double, N*N );
TESTING_MALLOC_DEV( ddA2, double, N*N );
TESTING_MALLOC_DEV( dtau2, double, min_mn );
TESTING_MALLOC_DEV( dwork, double, max(5*min_mn, (BLOCK_SIZE*2+2)*min_mn) );
TESTING_MALLOC_DEV( dwork2, double, max(5*min_mn, (BLOCK_SIZE*2+2)*min_mn) );
// todo replace with magma_dlaset
magmablas_dlaset( MagmaFull, N, N, c_zero, c_zero, ddA, N );
magmablas_dlaset( MagmaFull, N, N, c_zero, c_zero, d_T, N );
magmablas_dlaset( MagmaFull, N, N, c_zero, c_zero, ddA2, N );
magmablas_dlaset( MagmaFull, N, N, c_zero, c_zero, d_T2, N );
lwork = -1;
lapackf77_dgeqrf(&M, &N, NULL, &M, NULL, tmp, &lwork, &info);
lwork = (magma_int_t)MAGMA_D_REAL( tmp[0] );
lwork = max(lwork, N*N);
TESTING_MALLOC_CPU( h_work, double, lwork );
/* Initialize the matrix */
lapackf77_dlarnv( &ione, ISEED, &n2, h_A );
lapackf77_dlacpy( MagmaUpperLowerStr, &M, &N, h_A, &lda, h_R, &lda );
magma_dsetmatrix( M, N, h_R, lda, d_A, ldda );
magma_dsetmatrix( M, N, h_R, lda, d_A2, ldda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
gpu_time = magma_sync_wtime(0);
if (opts.version == 1)
magma_dgeqr2x_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info);
else if (opts.version == 2)
magma_dgeqr2x2_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info);
else if (opts.version == 3)
magma_dgeqr2x3_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info);
else {
printf( "call magma_dgeqr2x4_gpu\n" );
/*
Going through NULL stream is faster
Going through any stream is slower
Doing two streams in parallel is slower than doing them sequentially
Queuing happens on the NULL stream - user defined buffers are smaller?
*/
magma_dgeqr2x4_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, NULL, &info);
//magma_dgeqr2x4_gpu(M, N, d_A, ldda, dtau, d_T, ddA, dwork, &info, stream[1]);
//magma_dgeqr2x4_gpu(M, N, d_A2, ldda, dtau2, d_T2, ddA2, dwork2, &info, stream[0]);
//magma_dgeqr2x4_gpu(M, N, d_A2, ldda, dtau2, d_T2, ddA2, dwork2, &info, NULL);
//gflops *= 2;
}
gpu_time = magma_sync_wtime(0) - gpu_time;
gpu_perf = gflops / gpu_time;
if (info != 0) {
printf("magma_dgeqr2x_gpu version %d returned error %d: %s.\n",
(int) opts.version, (int) info, magma_strerror( info ));
}
else {
if ( opts.check ) {
/* =====================================================================
Check the result, following zqrt01 except using the reduced Q.
This works for any M,N (square, tall, wide).
=================================================================== */
magma_dgetmatrix( M, N, d_A, ldda, h_R, M );
magma_dgetmatrix( N, N, ddA, N, h_T, N );
magma_dgetmatrix( min_mn, 1, dtau, min_mn, tau, min_mn );
// Restore the upper triangular part of A before the check
for(int col=0; col < N; col++){
for(int row=0; row <= col; row++)
h_R[row + col*M] = h_T[row + col*N];
}
magma_int_t ldq = M;
magma_int_t ldr = min_mn;
double *Q, *R;
double *work;
TESTING_MALLOC_CPU( Q, double, ldq*min_mn ); // M by K
TESTING_MALLOC_CPU( R, double, ldr*N ); // K by N
TESTING_MALLOC_CPU( work, double, min_mn );
// generate M by K matrix Q, where K = min(M,N)
lapackf77_dlacpy( "Lower", &M, &min_mn, h_R, &M, Q, &ldq );
lapackf77_dorgqr( &M, &min_mn, &min_mn, Q, &ldq, tau, h_work, &lwork, &info );
assert( info == 0 );
// copy K by N matrix R
lapackf77_dlaset( "Lower", &min_mn, &N, &c_zero, &c_zero, R, &ldr );
lapackf77_dlacpy( "Upper", &min_mn, &N, h_R, &M, R, &ldr );
// error = || R - Q^H*A || / (N * ||A||)
blasf77_dgemm( "Conj", "NoTrans", &min_mn, &N, &M,
&c_neg_one, Q, &ldq, h_A, &lda, &c_one, R, &ldr );
double Anorm = lapackf77_dlange( "1", &M, &N, h_A, &lda, work );
error2 = lapackf77_dlange( "1", &min_mn, &N, R, &ldr, work );
if ( N > 0 && Anorm > 0 )
error2 /= (N*Anorm);
TESTING_FREE_CPU( Q ); Q = NULL;
TESTING_FREE_CPU( R ); R = NULL;
TESTING_FREE_CPU( work ); work = NULL;
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
cpu_time = magma_wtime();
//lapackf77_dgeqrf(&M, &N, h_A, &lda, tau, h_work, &lwork, &info);
lapackf77_dlacpy( MagmaUpperLowerStr, &M, &N, h_R, &M, h_A, &lda );
lapackf77_dlarft( MagmaForwardStr, MagmaColumnwiseStr,
&M, &N, h_A, &lda, tau, h_work, &N);
//magma_dgeqr2(&M, &N, h_A, &lda, tau, h_work, &info);
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
if (info != 0)
printf("lapackf77_dgeqrf returned error %d: %s.\n",
(int) info, magma_strerror( info ));
/* =====================================================================
Check the result compared to LAPACK
=================================================================== */
// Restore the upper triangular part of A before the check
for(int col=0; col < N; col++){
for(int row=0; row <= col; row++)
h_R[row + col*M] = h_T[row + col*N];
}
error = lapackf77_dlange("M", &M, &N, h_A, &lda, work);
blasf77_daxpy(&n2, &c_neg_one, h_A, &ione, h_R, &ione);
error = lapackf77_dlange("M", &M, &N, h_R, &lda, work) / (N * error);
// Check if T is the same
magma_dgetmatrix( N, N, d_T, N, h_T, N );
double terr = 0.;
for(int col=0; col < N; col++)
for(int row=0; row <= col; row++)
terr += ( MAGMA_D_ABS(h_work[row + col*N] - h_T[row + col*N])*
MAGMA_D_ABS(h_work[row + col*N] - h_T[row + col*N]) );
terr = sqrt( terr );
// If comparison to LAPACK fail, check || R - Q^H*A || / (N * ||A||)
// and print fail if both fails, otherwise print ok (*)
printf("%5d %5d %7.2f (%7.2f) %7.2f (%7.2f) %8.2e %8.2e %s\n",
(int) M, (int) N, cpu_perf, 1000.*cpu_time, gpu_perf, 1000.*gpu_time,
error2, terr, (error2 < tol ? "ok" : "failed" ));
status += ! (error2 < tol);
}
else {
printf("%5d %5d --- ( --- ) %7.2f (%7.2f) --- \n",
(int) M, (int) N, gpu_perf, 1000.*gpu_time);
}
}
TESTING_FREE_CPU( tau );
TESTING_FREE_CPU( h_A );
TESTING_FREE_CPU( h_T );
TESTING_FREE_CPU( h_work );
TESTING_FREE_PIN( h_R );
TESTING_FREE_DEV( d_A );
TESTING_FREE_DEV( d_T );
TESTING_FREE_DEV( ddA );
TESTING_FREE_DEV( dtau );
TESTING_FREE_DEV( dwork );
TESTING_FREE_DEV( d_A2 );
TESTING_FREE_DEV( d_T2 );
TESTING_FREE_DEV( ddA2 );
TESTING_FREE_DEV( dtau2 );
TESTING_FREE_DEV( dwork2 );
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
TESTING_FINALIZE();
return status;
}
/**
Purpose
-------
DORGQR generates an M-by-N DOUBLE_PRECISION matrix Q with orthonormal columns,
which is defined as the first N columns of a product of K elementary
reflectors of order M
Q = H(1) H(2) . . . H(k)
as returned by DGEQRF_GPU.
Arguments
---------
@param[in]
m INTEGER
The number of rows of the matrix Q. M >= 0.
@param[in]
n INTEGER
The number of columns of the matrix Q. M >= N >= 0.
@param[in]
k INTEGER
The number of elementary reflectors whose product defines the
matrix Q. N >= K >= 0.
@param[in,out]
dA DOUBLE_PRECISION array A on the GPU, dimension (LDDA,N).
On entry, the i-th column must contain the vector
which defines the elementary reflector H(i), for
i = 1,2,...,k, as returned by DGEQRF_GPU in the
first k columns of its array argument A.
On exit, the M-by-N matrix Q.
@param[in]
ldda INTEGER
The first dimension of the array A. LDDA >= max(1,M).
@param[in]
tau DOUBLE_PRECISION array, dimension (K)
TAU(i) must contain the scalar factor of the elementary
reflector H(i), as returned by DGEQRF_GPU.
@param[in]
dT (workspace) DOUBLE_PRECISION work space array on the GPU,
dimension (2*MIN(M, N) + (N+31)/32*32 )*NB.
This must be the 6th argument of magma_dgeqrf_gpu
[ note that if N here is bigger than N in magma_dgeqrf_gpu,
the workspace requirement DT in magma_dgeqrf_gpu must be
as specified in this routine ].
@param[in]
nb INTEGER
This is the block size used in DGEQRF_GPU, and correspondingly
the size of the T matrices, used in the factorization, and
stored in DT.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument has an illegal value
@ingroup magma_dgeqrf_comp
********************************************************************/
extern "C" magma_int_t
magma_dorgqr_gpu(
magma_int_t m, magma_int_t n, magma_int_t k,
magmaDouble_ptr dA, magma_int_t ldda,
double *tau,
magmaDouble_ptr dT, magma_int_t nb,
magma_int_t *info)
{
#define dA(i,j) (dA + (i) + (j)*ldda)
#define dT(j) (dT + (j)*nb)
double c_zero = MAGMA_D_ZERO;
double c_one = MAGMA_D_ONE;
magma_int_t m_kk, n_kk, k_kk, mi;
magma_int_t lwork, lpanel;
magma_int_t i, ib, ki, kk, iinfo;
magma_int_t lddwork;
magmaDouble_ptr dV, dW;
double *work, *panel;
*info = 0;
if (m < 0) {
*info = -1;
} else if ((n < 0) || (n > m)) {
*info = -2;
} else if ((k < 0) || (k > n)) {
*info = -3;
} else if (ldda < max(1,m)) {
*info = -5;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if (n <= 0) {
return *info;
}
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
// first kk columns are handled by blocked method.
// ki is start of 2nd-to-last block
if ((nb > 1) && (nb < k)) {
ki = (k - nb - 1) / nb * nb;
kk = min( k, ki+nb );
} else {
ki = 0;
kk = 0;
}
// Allocate CPU work space
// n*nb for dorgqr workspace
// (m - kk)*(n - kk) for last block's panel
lwork = n*nb;
lpanel = (m - kk)*(n - kk);
magma_dmalloc_cpu( &work, lwork + lpanel );
if ( work == NULL ) {
*info = MAGMA_ERR_HOST_ALLOC;
return *info;
}
panel = work + lwork;
// Allocate work space on GPU
if (MAGMA_SUCCESS != magma_dmalloc( &dV, ldda*nb )) {
magma_free_cpu( work );
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
// dT workspace has:
// 2*min(m,n)*nb for T and R^{-1} matrices from geqrf
// ((n+31)/32*32 )*nb for dW larfb workspace.
lddwork = min(m,n);
dW = dT + 2*lddwork*nb;
magma_queue_t stream;
magma_queue_create( &stream );
// Use unblocked code for the last or only block.
if (kk < n) {
m_kk = m - kk;
n_kk = n - kk;
k_kk = k - kk;
magma_dgetmatrix( m_kk, k_kk,
dA(kk, kk), ldda, panel, m_kk );
lapackf77_dorgqr( &m_kk, &n_kk, &k_kk,
panel, &m_kk,
&tau[kk], work, &lwork, &iinfo );
magma_dsetmatrix( m_kk, n_kk,
panel, m_kk, dA(kk, kk), ldda );
// Set A(1:kk,kk+1:n) to zero.
magmablas_dlaset( MagmaFull, kk, n - kk, c_zero, c_zero, dA(0, kk), ldda );
}
if (kk > 0) {
// Use blocked code
// stream: copy Aii to V --> laset --> laset --> larfb --> [next]
// CPU has no computation
magmablasSetKernelStream( stream );
for (i = ki; i >= 0; i -= nb) {
ib = min( nb, k-i );
mi = m - i;
// Copy current panel on the GPU from dA to dV
magma_dcopymatrix_async( mi, ib,
dA(i,i), ldda,
dV, ldda, stream );
// set panel to identity
magmablas_dlaset( MagmaFull, i, ib, c_zero, c_zero, dA(0, i), ldda );
magmablas_dlaset( MagmaFull, mi, ib, c_zero, c_one, dA(i, i), ldda );
if (i < n) {
// Apply H to A(i:m,i:n) from the left
magma_dlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise,
mi, n-i, ib,
dV, ldda, dT(i), nb,
dA(i, i), ldda, dW, lddwork );
}
}
}
magma_queue_sync( stream );
magma_free( dV );
magma_free_cpu( work );
magma_queue_destroy( stream );
magmablasSetKernelStream( orig_stream );
return *info;
} /* magma_dorgqr_gpu */
/***************************************************************************//**
Purpose
-------
DORGBR generates one of the real orthogonal matrices Q or P**H
determined by DGEBRD when reducing a real matrix A to bidiagonal
form: A = Q * B * P**H. Q and P**H are defined as products of
elementary reflectors H(i) or G(i) respectively.
If VECT = MagmaQ, A is assumed to have been an M-by-K matrix, and Q
is of order M:
if m >= k, Q = H(1) H(2) . . . H(k) and DORGBR returns the first n
columns of Q, where m >= n >= k;
if m < k, Q = H(1) H(2) . . . H(m-1) and DORGBR returns Q as an
M-by-M matrix.
If VECT = MagmaP, A is assumed to have been a K-by-N matrix, and P**H
is of order N:
if k < n, P**H = G(k) . . . G(2) G(1) and DORGBR returns the first m
rows of P**H, where n >= m >= k;
if k >= n, P**H = G(n-1) . . . G(2) G(1) and DORGBR returns P**H as
an N-by-N matrix.
Arguments
---------
@param[in]
vect magma_vect_t
Specifies whether the matrix Q or the matrix P**H is
required, as defined in the transformation applied by DGEBRD:
= MagmaQ: generate Q;
= MagmaP: generate P**H.
@param[in]
m magma_int_t
The number of rows of the matrix Q or P**H to be returned.
M >= 0.
@param[in]
n magma_int_t
The number of columns of the matrix Q or P**H to be returned.
N >= 0.
If VECT = MagmaQ, M >= N >= min(M,K);
if VECT = MagmaP, N >= M >= min(N,K).
@param[in]
k magma_int_t
If VECT = MagmaQ, the number of columns in the original M-by-K
matrix reduced by DGEBRD.
If VECT = MagmaP, the number of rows in the original K-by-N
matrix reduced by DGEBRD.
K >= 0.
@param[in,out]
A double array, dimension (LDA,N)
On entry, the vectors which define the elementary reflectors,
as returned by DGEBRD.
On exit, the M-by-N matrix Q or P**H.
@param[in]
lda magma_int_t
The leading dimension of the array A. LDA >= M.
@param[in]
tau double array, dimension
(min(M,K)) if VECT = MagmaQ
(min(N,K)) if VECT = MagmaP
TAU(i) must contain the scalar factor of the elementary
reflector H(i) or G(i), which determines Q or P**H, as
returned by DGEBRD in its array argument TAUQ or TAUP.
@param[out]
work double array, dimension (MAX(1,LWORK))
On exit, if *info = 0, WORK(1) returns the optimal LWORK.
@param[in]
lwork magma_int_t
The dimension of the array WORK. LWORK >= max(1,min(M,N)).
For optimum performance LWORK >= min(M,N)*NB, where NB
is the optimal blocksize.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
info magma_int_t
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_ungbr
*******************************************************************************/
extern "C" magma_int_t
magma_dorgbr(
magma_vect_t vect, magma_int_t m, magma_int_t n, magma_int_t k,
double *A, magma_int_t lda,
double *tau,
double *work, magma_int_t lwork,
magma_int_t *info)
{
#define A(i_,j_) (A + (i_) + (j_)*lda)
// Constants
const double c_zero = MAGMA_D_ZERO;
const double c_one = MAGMA_D_ONE;
const magma_int_t ineg_one = -1;
// Local variables
bool lquery, wantq;
magma_int_t i, iinfo, j, lwkmin, lwkopt, min_mn;
// Test the input arguments
*info = 0;
wantq = (vect == MagmaQ);
min_mn = min( m, n );
lquery = (lwork == -1);
if ( ! wantq && vect != MagmaP ) {
*info = -1;
} else if (m < 0) {
*info = -2;
} else if (n < 0 || (wantq && (n > m || n < min(m,k))) || ( ! wantq && (m > n || m < min(n,k)))) {
*info = -3;
} else if (k < 0) {
*info = -4;
} else if (lda < max( 1, m )) {
*info = -6;
}
// Check workspace size
if (*info == 0) {
work[0] = c_one;
if (wantq) {
if (m >= k) {
// magma_dorgqr takes dT instead of work
// magma_dorgqr2 doesn't take work
//magma_dorgqr2( m, n, k, A, lda, tau, work, -1, &iinfo );
lapackf77_dorgqr( &m, &n, &k, A, &lda, tau, work, &ineg_one, &iinfo );
}
else if (m > 1) {
//magma_dorgqr2( m-1, m-1, m-1, A(1,1), lda, tau, work, -1, &iinfo );
magma_int_t m1 = m-1;
lapackf77_dorgqr( &m1, &m1, &m1, A(1,1), &lda, tau, work, &ineg_one, &iinfo );
}
lwkopt = MAGMA_D_REAL( work[0] );
lwkmin = min_mn;
}
else {
if (k < n) {
magma_dorglq( m, n, k, A, lda, tau, work, -1, &iinfo );
}
else if (n > 1) {
magma_dorglq( n-1, n-1, n-1, A(1,1), lda, tau, work, -1, &iinfo );
}
lwkopt = MAGMA_D_REAL( work[0] );
lwkmin = lwkopt;
}
if (lwork < lwkmin && ! lquery) {
*info = -9;
}
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
else if (lquery) {
work[0] = magma_dmake_lwork( lwkopt );
return *info;
}
// Quick return if possible
if (m == 0 || n == 0) {
work[0] = c_one;
return *info;
}
if (wantq) {
// Form Q, determined by a call to DGEBRD to reduce an m-by-k
// matrix
if (m >= k) {
// If m >= k, assume m >= n >= k
magma_dorgqr2( m, n, k, A, lda, tau, /*work, lwork,*/ &iinfo );
}
else {
// If m < k, assume m = n
// Shift the vectors which define the elementary reflectors one
// column to the right, and set the first row and column of Q
// to those of the unit matrix
for (j=m-1; j >= 1; --j) {
*A(0,j) = c_zero;
for (i=j + 1; i < m; ++i) {
*A(i,j) = *A(i,j-1);
}
}
*A(0,0) = c_one;
for (i=1; i < m; ++i) {
*A(i,0) = c_zero;
}
if (m > 1) {
// Form Q(2:m,2:m)
magma_dorgqr2( m-1, m-1, m-1, A(1,1), lda, tau, /*work, lwork,*/ &iinfo );
}
}
}
else {
// Form P**H, determined by a call to DGEBRD to reduce a k-by-n
// matrix
if (k < n) {
// If k < n, assume k <= m <= n
magma_dorglq( m, n, k, A, lda, tau, work, lwork, &iinfo );
}
else {
// If k >= n, assume m = n
// Shift the vectors which define the elementary reflectors one
// row downward, and set the first row and column of P**H to
// those of the unit matrix
*A(0,0) = c_one;
for (i=1; i < n; ++i) {
*A(i,0) = c_zero;
}
for (j=1; j < n; ++j) {
for (i=j-1; i >= 1; --i) {
*A(i,j) = *A(i-1,j);
}
*A(0,j) = c_zero;
}
if (n > 1) {
// Form P**H(2:n,2:n)
magma_dorglq( n-1, n-1, n-1, A(1,1), lda, tau, work, lwork, &iinfo );
}
}
}
work[0] = magma_dmake_lwork( lwkopt );
return *info;
}
