//===================================================================================================================
//===================================================================================================================
//===================================================================================================================
extern "C" void
magma_slarft_sm32x32_batched(magma_int_t n, magma_int_t k,
float **v_array, magma_int_t ldv,
float **tau_array,
float **T_array, magma_int_t ldt,
magma_int_t batchCount, magma_queue_t queue)
{
if ( k <= 0) return;
//==================================
// GEMV
//==================================
#define USE_GEMV2
#define use_gemm_larft_sm32
#if defined(use_gemm_larft_sm32)
magma_sgemm_batched( MagmaConjTrans, MagmaNoTrans,
k, k, n,
MAGMA_S_ONE, v_array, ldv,
v_array, ldv,
MAGMA_S_ZERO, T_array, ldt,
batchCount, queue );
magmablas_slaset_batched( MagmaLower, k, k,
MAGMA_S_ZERO, MAGMA_S_ZERO,
T_array, ldt, batchCount, queue );
#else
#if 1
for (magma_int_t i=0; i < k; i++)
{
//W(1:i-1) := - tau(i) * V(i:n,1:i-1)' * V(i:n,i)
//T( i, i ) = tau( i )
//custom implementation.
#ifdef USE_GEMV2
magmablas_slarft_gemvrowwise_batched( n-i, i,
tau_array,
v_array, ldv,
T_array, ldt,
batchCount, queue);
#else
magmablas_slarft_gemvcolwise_batched( n-i, i, v_array, ldv, T_array, ldt, tau_array, batchCount, queue);
#endif
}
#else
//seems to be very slow when k=32 while the one by one loop above is faster
slarft_gemv_loop_inside_kernel_batched(n, k, tau_array, v_array, ldv, T_array, ldt, batchCount, queue);
#endif
#endif
//==================================
// TRMV
//==================================
//T(1:i-1,i) := T(1:i-1,1:i-1) * W(1:i-1) i=[1:k]
magmablas_slarft_strmv_sm32x32_batched(k, k, tau_array, T_array, ldt, T_array, ldt, batchCount, queue);
}
extern "C" magma_int_t
magma_sgetf2_nopiv_batched(
magma_int_t m, magma_int_t n,
float **dA_array, magma_int_t ldda,
float **dW0_displ,
float **dW1_displ,
float **dW2_displ,
magma_int_t *info_array,
magma_int_t gbstep,
magma_int_t batchCount, magma_queue_t queue)
{
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) {
return arginfo;
}
float c_neg_one = MAGMA_S_NEG_ONE;
float c_one = MAGMA_S_ONE;
magma_int_t nb = BATF2_NB;
magma_int_t min_mn = min(m, n);
magma_int_t gbj, panelj, step, ib;
for( panelj=0; panelj < min_mn; panelj += nb)
{
ib = min(nb, min_mn-panelj);
for (step=0; step < ib; step++) {
gbj = panelj+step;
#if 0
size_t required_shmem_size = ((m-panelj)*ib)*sizeof(float);
if ( required_shmem_size > (MAX_SHARED_ALLOWED*1024))
#else
if ( (m-panelj) > 0)
#endif
{
// Compute elements J+1:M of J-th column.
if (gbj < m) {
arginfo = magma_sscal_sger_batched( m-gbj, ib-step, gbj, dA_array, ldda, info_array, gbstep, batchCount, queue );
if (arginfo != 0 ) return arginfo;
}
}
else {
// TODO
}
}
if ( (n-panelj-ib) > 0) {
// continue the update of the selected ib row column panelj+ib:n(TRSM)
magma_sgetf2trsm_batched(ib, n-panelj-ib, dA_array, panelj, ldda, batchCount, queue);
// do the blocked DGER = DGEMM for the remaining panelj+ib:n columns
magma_sdisplace_pointers(dW0_displ, dA_array, ldda, ib+panelj, panelj, batchCount, queue);
magma_sdisplace_pointers(dW1_displ, dA_array, ldda, panelj, ib+panelj, batchCount, queue);
magma_sdisplace_pointers(dW2_displ, dA_array, ldda, ib+panelj, ib+panelj, batchCount, queue);
magma_sgemm_batched( MagmaNoTrans, MagmaNoTrans, m-(panelj+ib), n-(panelj+ib), ib,
c_neg_one, dW0_displ, ldda,
dW1_displ, ldda,
c_one, dW2_displ, ldda,
batchCount, queue );
}
}
//magma_free_cpu(cpuAarray);
return 0;
}
/***************************************************************************//**
Purpose
-------
This is an internal routine that might have many assumption.
Documentation is not fully completed
SGETRF_PANEL 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]
min_recpnb INTEGER.
Internal use. The recursive nb
@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]
dipiv_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]
dpivinfo_array Array of pointers, dimension (batchCount), for internal use.
@param[in,out]
dX_array Array of pointers, dimension (batchCount).
Each is a REAL array X of dimension ( lddx, n ).
On entry, should be set to 0
On exit, the solution matrix X
@param[in]
dX_length INTEGER.
The size of each workspace matrix dX
@param[in,out]
dinvA_array Array of pointers, dimension (batchCount).
Each is a REAL array dinvA, a workspace on device.
If side == MagmaLeft, dinvA must be of size >= ceil(m/STRTRI_BATCHED_NB)*STRTRI_BATCHED_NB*STRTRI_BATCHED_NB,
If side == MagmaRight, dinvA must be of size >= ceil(n/STRTRI_BATCHED_NB)*STRTRI_BATCHED_NB*STRTRI_BATCHED_NB,
@param[in]
dinvA_length INTEGER
The size of each workspace matrix dinvA
@param[in]
dW1_displ Workspace array of pointers, for internal use.
@param[in]
dW2_displ Workspace array of pointers, for internal use.
@param[in]
dW3_displ Workspace array of pointers, for internal use.
@param[in]
dW4_displ Workspace array of pointers, for internal use.
@param[in]
dW5_displ Workspace array of pointers, for internal use.
@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]
gbstep INTEGER
internal use.
@param[in]
batchCount INTEGER
The number of matrices to operate on.
@param[in]
queue magma_queue_t
Queue to execute in.
@ingroup magma_getrf_batched
*******************************************************************************/
extern "C" magma_int_t
magma_sgetrf_recpanel_batched(
magma_int_t m, magma_int_t n, magma_int_t min_recpnb,
float** dA_array, magma_int_t ldda,
magma_int_t** dipiv_array, magma_int_t** dpivinfo_array,
float** dX_array, magma_int_t dX_length,
float** dinvA_array, magma_int_t dinvA_length,
float** dW1_displ, float** dW2_displ,
float** dW3_displ, float** dW4_displ,
float** dW5_displ,
magma_int_t *info_array, magma_int_t gbstep,
magma_int_t batchCount, magma_queue_t queue)
{
//magma_int_t DEBUG = 3;
// Quick return if possible
if (m == 0 || n == 0) {
return 0;
}
float **dA_displ = NULL;
magma_malloc((void**)&dA_displ, batchCount * sizeof(*dA_displ));
magma_int_t **dipiv_displ = NULL;
magma_malloc((void**)&dipiv_displ, batchCount * sizeof(*dipiv_displ));
magma_int_t panel_nb = n;
if (panel_nb <= min_recpnb) {
//if (DEBUG > 0)printf("calling bottom panel recursive with m=%d nb=%d\n",m,n);
// panel factorization
//magma_sdisplace_pointers(dA_displ, dA_array, ldda, 0, 0, batchCount);
magma_sgetf2_batched(m, panel_nb,
dA_array, ldda,
dW1_displ, dW2_displ, dW3_displ,
dipiv_array, info_array, gbstep, batchCount, queue);
}
else {
// split A over two [A A2]
// panel on A1, update on A2 then panel on A1
magma_int_t n1 = n/2;
magma_int_t n2 = n-n1;
magma_int_t m1 = m;
magma_int_t m2 = m-n1;
magma_int_t p1 = 0;
magma_int_t p2 = n1;
// panel on A1
//if (DEBUG > 0)printf("calling recursive panel on A1 with m=%d nb=%d min_recpnb %d\n",m1,n1,min_recpnb);
magma_sdisplace_pointers(dA_displ, dA_array, ldda, p1, p1, batchCount, queue);
magma_idisplace_pointers(dipiv_displ, dipiv_array, 1, p1, 0, batchCount, queue);
magma_sgetrf_recpanel_batched(
m1, n1, min_recpnb,
dA_displ, ldda,
dipiv_displ, dpivinfo_array,
dX_array, dX_length,
dinvA_array, dinvA_length,
dW1_displ, dW2_displ,
dW3_displ, dW4_displ, dW5_displ,
info_array, gbstep, batchCount, queue);
// update A2
//if (DEBUG > 0)printf("calling TRSM with m=%d n=%d\n",m1,n2);
// setup pivinfo
setup_pivinfo_batched(dpivinfo_array, dipiv_displ, m1, n1, batchCount, queue);
magma_sdisplace_pointers(dW5_displ, dA_array, ldda, p1, p2, batchCount, queue);
magma_slaswp_rowparallel_batched( n2, dW5_displ, ldda,
dX_array, n1,
0, n1,
dpivinfo_array, batchCount, queue );
magmablas_strsm_outofplace_batched( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit, 1,
n1, n2,
MAGMA_S_ONE,
dA_displ, ldda, // dA
dX_array, n1, // dB
dW5_displ, ldda, // dX
dinvA_array, dinvA_length,
dW1_displ, dW2_displ,
dW3_displ, dW4_displ,
0, batchCount, queue );
magma_sdisplace_pointers(dW1_displ, dA_array, ldda, p2, 0, batchCount, queue);
magma_sdisplace_pointers(dA_displ, dA_array, ldda, p2, p2, batchCount, queue);
//if (DEBUG > 0)printf("calling update A2(%d,%d) -= A(%d,%d)*A(%d,%d) with m=%d n=%d k=%d ldda %d\n",p2,p2,p2,0,p1,p2,m2,n2,n1,ldda);
magma_sgemm_batched( MagmaNoTrans, MagmaNoTrans, m2, n2, n1,
MAGMA_S_NEG_ONE, dW1_displ, ldda,
dW5_displ, ldda,
MAGMA_S_ONE, dA_displ, ldda,
batchCount, queue );
// panel on A2
//if (DEBUG > 0)printf("calling recursive panel on A2 with m=%d nb=%d min_recpnb %d\n",m2,n2,min_recpnb);
magma_idisplace_pointers(dipiv_displ, dipiv_array, 1, p2, 0, batchCount, queue);
magma_sgetrf_recpanel_batched(
m2, n2, min_recpnb,
dA_displ, ldda,
dipiv_displ, dpivinfo_array,
dX_array, dX_length,
dinvA_array, dinvA_length,
dW1_displ, dW2_displ,
dW3_displ, dW4_displ, dW5_displ,
info_array, gbstep+p2, batchCount, queue);
// setup pivinfo
setup_pivinfo_batched(dpivinfo_array, dipiv_displ, m2, n2, batchCount, queue);
adjust_ipiv_batched(dipiv_displ, n2, n1, batchCount, queue);
magma_sdisplace_pointers(dW1_displ, dA_array, ldda, p2, 0, batchCount, queue); // no need since it is above
magma_slaswp_rowparallel_batched( n1, dW1_displ, ldda,
dW1_displ, ldda,
n1, n,
dpivinfo_array, batchCount, queue );
}
magma_free(dA_displ);
magma_free(dipiv_displ);
return 0;
}
/**
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;
}
/**
Purpose
-------
SLARFB applies a real block reflector H or its transpose H^H to a
REAL m by n matrix C, from the left.
__Note that this function assumes__ that the upper part of dV_array is 0
because it is referenced. Same for upper/lower part of dT_array.
Arguments
---------
@param[in]
side magma_side_t
- = MagmaLeft: apply H or H^H from the Left
- = MagmaRight: apply H or H^H from the Right
@param[in]
trans magma_trans_t
- = MagmaNoTrans: apply H (No transpose)
- = MagmaTrans: apply H^H (Conjugate transpose)
@param[in]
direct magma_direct_t
Indicates how H is formed from a product of elementary
reflectors
- = MagmaForward: H = H(1) H(2) . . . H(k) (Forward)
- = MagmaBackward: H = H(k) . . . H(2) H(1) (Backward)
@param[in]
storev magma_storev_t
Indicates how the vectors which define the elementary
reflectors are stored:
- = MagmaColumnwise: Columnwise
- = MagmaRowwise: Rowwise
@param[in]
m INTEGER
The number of rows of the matrix C.
@param[in]
n INTEGER
The number of columns of the matrix C.
@param[in]
k INTEGER
The order of the matrix T (= the number of elementary
reflectors whose product defines the block reflector).
@param[in]
dV_array REAL array on the GPU, dimension
(LDDV,K) if STOREV = MagmaColumnwise
(LDDV,M) if STOREV = MagmaRowwise and SIDE = MagmaLeft
(LDDV,N) if STOREV = MagmaRowwise and SIDE = MagmaRight
The matrix V. See further details.
@param[in]
lddv INTEGER
The leading dimension of the array V.
If STOREV = MagmaColumnwise and SIDE = MagmaLeft, LDDV >= max(1,M);
if STOREV = MagmaColumnwise and SIDE = MagmaRight, LDDV >= max(1,N);
if STOREV = MagmaRowwise, LDDV >= K.
@param[in]
dT_array REAL array on the GPU, dimension (LDDT,K)
The triangular k by k matrix T in the representation of the
block reflector.
@param[in]
lddt INTEGER
The leading dimension of the array T. LDDT >= K.
@param[in,out]
dC_array REAL array on the GPU, dimension (LDDC,N)
On entry, the m by n matrix C.
On exit, C is overwritten by H*C, or H^H*C, or C*H, or C*H^H.
@param[in]
lddc INTEGER
The leading dimension of the array C. LDA >= max(1,M).
@param
dwork_array (workspace) REAL array, dimension (LDWORK,K)
@param[in]
ldwork INTEGER
The leading dimension of the array WORK.
If SIDE = MagmaLeft, LDWORK >= max(1,N);
if SIDE = MagmaRight, LDWORK >= max(1,M);
@param
dworkvt_array (workspace) REAL array, dimension (LDWORKT,K)
@param[in]
ldworkvt INTEGER
The leading dimension of the array WORKVT.
LDWORKVT >= max(1,min(M,N));
@param[in]
batchCount INTEGER
The number of matrices to operate on.
@param[in]
queue magma_queue_t
Queue to execute in.
Further Details
---------------
The shape of the matrix V and the storage of the vectors which define
the H(i) is best illustrated by the following example with n = 5 and
k = 3.
All elements including 0's and 1's are stored, unlike LAPACK.
DIRECT = MagmaForward and DIRECT = MagmaForward and
STOREV = MagmaColumnwise: STOREV = MagmaRowwise:
V = ( 1 0 0 ) V = ( 1 v1 v1 v1 v1 )
( v1 1 0 ) ( 0 1 v2 v2 v2 )
( v1 v2 1 ) ( 0 0 1 v3 v3 )
( v1 v2 v3 )
( v1 v2 v3 )
DIRECT = MagmaBackward and DIRECT = MagmaBackward and
STOREV = MagmaColumnwise: STOREV = MagmaRowwise:
V = ( v1 v2 v3 ) V = ( v1 v1 1 0 0 )
( v1 v2 v3 ) ( v2 v2 v2 1 0 )
( 1 v2 v3 ) ( v3 v3 v3 v3 1 )
( 0 1 v3 )
( 0 0 1 )
@ingroup magma_saux3
********************************************************************/
extern "C" magma_int_t
magma_slarfb_gemm_batched(
magma_side_t side, magma_trans_t trans, magma_direct_t direct, magma_storev_t storev,
magma_int_t m, magma_int_t n, magma_int_t k,
magmaFloat_const_ptr dV_array[], magma_int_t lddv,
magmaFloat_const_ptr dT_array[], magma_int_t lddt,
magmaFloat_ptr dC_array[], magma_int_t lddc,
magmaFloat_ptr dwork_array[], magma_int_t ldwork,
magmaFloat_ptr dworkvt_array[], magma_int_t ldworkvt,
magma_int_t batchCount, magma_queue_t queue)
{
float c_zero = MAGMA_S_ZERO;
float c_one = MAGMA_S_ONE;
float c_neg_one = MAGMA_S_NEG_ONE;
/* Function Body */
magma_int_t info = 0;
if (m <= 0 || n <= 0) {
return info;
}
// internal variable
magma_int_t ldwvt = (m > n ? k : m);
magma_int_t ldw;
if ( side == MagmaLeft ) {
ldw = k;
} else {
ldw = m;
}
// opposite of trans
magma_trans_t transt;
if (trans == MagmaNoTrans)
transt = MagmaTrans;
else
transt = MagmaNoTrans;
MAGMA_UNUSED( transt ); // TODO: is this a bug that it isn't used?
// whether V is stored transposed or not
magma_trans_t notransV, transV;
if (storev == MagmaColumnwise) {
notransV = MagmaNoTrans;
transV = MagmaTrans;
}
else {
notransV = MagmaTrans;
transV = MagmaNoTrans;
}
if ( side == MagmaLeft ) {
// Form H C or H^H C
// Comments assume H C.
// When forming H^H C, T gets transposed via transt for m >= n or by trans for m < n.
// W = V' C
magma_sgemm_batched( MagmaTrans,notransV, /*NontransLeft*/
k, n, m,
c_one, dV_array, lddv,
dC_array, lddc,
c_zero, dwork_array, ldw,
batchCount, queue );
if (m <= n) {
// W2 = V T
magma_sgemm_batched( notransV, trans, /* (NoTrans), trans(ConjTrans),*/
m, k, k,
c_one, dV_array, lddv,
dT_array, lddt,
c_zero, dworkvt_array, ldwvt,
batchCount, queue );
// C = C - W2 W = C - V T V' C = (I - V T V') C = H C
magma_sgemm_batched( MagmaNoTrans, MagmaNoTrans,
m, n, k,
c_neg_one, dworkvt_array, ldwvt,
dwork_array, ldw,
c_one, dC_array, lddc,
batchCount, queue );
}
else {
// W2 = T W = T V' C
magma_sgemm_batched( trans, MagmaNoTrans,
k, n, k,
c_one, dT_array, lddt,
dwork_array, ldw,
c_zero, dworkvt_array, ldwvt,
batchCount, queue );
// C = C - V W2 = C - V T V' C = (I - V T V') C = H C
magma_sgemm_batched( notransV, MagmaNoTrans,
m, n, k,
c_neg_one, dV_array, lddv,
dworkvt_array, ldwvt,
c_one, dC_array, lddc,
batchCount, queue );
}
}
else {
// Form C H or C H^H
// Comments assume C H.
// When forming C H^H, T gets transposed via trans.
// W = C V
magma_sgemm_batched( MagmaNoTrans, notransV,
m, k, n,
c_one, dC_array, lddc,
dV_array, lddv,
c_zero, dwork_array, ldw,
batchCount, queue );
if (m <= n) {
// W2 = W T = C V T
magma_sgemm_batched( MagmaNoTrans, trans,
m, k, k,
c_one, dwork_array, ldw,
dT_array, lddt,
c_zero, dworkvt_array, ldwvt,
batchCount, queue );
// C = C - W2 V' = C - C V T V' = C (I - V T V') = C H
magma_sgemm_batched( MagmaNoTrans, transV,
m, n, k,
c_neg_one, dworkvt_array, ldwvt,
dV_array, lddv,
c_one, dC_array, lddc,
batchCount, queue );
}
else {
// W2 = T V'
magma_sgemm_batched( trans, transV,
k, n, k,
c_one, dT_array, lddt,
dV_array, lddv,
c_zero, dworkvt_array, ldwvt,
batchCount, queue );
// C = C - W W2 = C - C V T V' = C (I - V T V') = C H
magma_sgemm_batched( MagmaNoTrans, MagmaNoTrans,
m, n, k,
c_neg_one, dwork_array, ldw,
dworkvt_array, ldwvt,
c_one, dC_array, lddc,
batchCount, queue );
}
}
return MAGMA_SUCCESS;
} /* magma_slarfb */
//===================================================================================================================
//===================================================================================================================
//===================================================================================================================
extern "C" magma_int_t
magma_slarft_batched(magma_int_t n, magma_int_t k, magma_int_t stair_T,
float **v_array, magma_int_t ldv,
float **tau_array, float **T_array, magma_int_t ldt,
float **work_array, magma_int_t lwork,
magma_int_t batchCount, magma_queue_t queue)
{
float c_one = MAGMA_S_ONE;
float c_zero = MAGMA_S_ZERO;
if ( k <= 0) return 0;
if ( stair_T > 0 && k <= stair_T) return 0;
magma_int_t maxnb = max_shared_bsiz;
if ( lwork < k*ldt)
{
magma_xerbla( __func__, -(10) );
return -10;
}
if ( stair_T > 0 && stair_T > maxnb)
{
magma_xerbla( __func__, -(3) );
return -3;
}
magma_int_t DEBUG=0;
magma_int_t nb = stair_T == 0 ? min(k,maxnb) : stair_T;
magma_int_t i, j, prev_n, mycol, rows;
float **dW1_displ = NULL;
float **dW2_displ = NULL;
float **dW3_displ = NULL;
float **dTstep_array = NULL;
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**)&dTstep_array, batchCount * sizeof(*dTstep_array));
//float *Tstep = k > nb ? work : T;
if (k > nb)
{
magma_sdisplace_pointers(dTstep_array, work_array, lwork, 0, 0, batchCount, queue);
}
else
{
magma_sdisplace_pointers(dTstep_array, T_array, ldt, 0, 0, batchCount, queue);
}
//magma_int_t ldtstep = k > nb ? k : ldt;
magma_int_t ldtstep = ldt; //a enlever
// stair_T = 0 meaning all T
// stair_T > 0 meaning the triangular portion of T has been computed.
// the value of stair_T is the nb of these triangulars
//GEMV compute the whole triangular upper portion of T (phase 1)
// TODO addcublas to check perf
magma_sgemm_batched( MagmaConjTrans, MagmaNoTrans,
k, k, n,
c_one, v_array, ldv,
v_array, ldv,
c_zero, dTstep_array, ldtstep,
batchCount, queue );
magmablas_slaset_batched( MagmaLower, k, k, MAGMA_S_ZERO, MAGMA_S_ZERO, dTstep_array, ldtstep, batchCount, queue );
// no need for it as T is expected to be lower zero
//if (k > nb) magmablas_slaset_batched( MagmaLower, k, k, MAGMA_S_ZERO, MAGMA_S_ZERO, dTstep_array, ldtstep, batchCount, queue );
//TRMV
//T(1:i-1,i) := T(1:i-1,1:i-1) * W(1:i-1) i=[1:k]
// TRMV is split over block of column of size nb
// the update should be done from top to bottom so:
// 1- a gemm using the previous computed columns
// of T to update rectangular upper protion above
// the triangle of my columns
// 2- the columns need to be updated by a serial
// loop over of gemv over itself. since we limit the
// shared memory to nb, this nb column
// are split vertically by chunk of nb rows
dim3 grid(1, 1, batchCount);
for (j=0; j < k; j += nb)
{
prev_n = j;
mycol = min(nb, k-j);
// note that myrow = prev_n + mycol;
if (prev_n > 0 && mycol > 0) {
if (DEBUG == 3) {
printf("doing gemm on the rectangular portion of size %d %d of T(%d,%d)\n",
(int) prev_n, (int) mycol, 0, (int) j );
}
magma_sdisplace_pointers(dW1_displ, dTstep_array, ldtstep, 0, j, batchCount, queue);
magma_sdisplace_pointers(dW2_displ, T_array, ldt, 0, j, batchCount, queue);
magma_sgemm_batched( MagmaNoTrans, MagmaNoTrans,
prev_n, mycol, prev_n,
c_one, T_array, ldt,
dW1_displ, ldtstep,
c_zero, dW2_displ, ldt,
batchCount, queue );
// update my rectangular portion (prev_n,mycol) using sequence of gemv
magma_sdisplace_pointers(dW1_displ, dTstep_array, ldtstep, j, j, batchCount, queue);
magma_sdisplace_pointers(dW3_displ, tau_array, 1, j, 0, batchCount, queue);
for (i=0; i < prev_n; i += nb)
{
rows = min(nb,prev_n-i);
if (DEBUG == 3) {
printf(" doing recstrmv on the rectangular portion of size %d %d of T(%d,%d)\n",
(int) rows, (int) mycol, (int) i, (int) j );
}
if (rows > 0 && mycol > 0)
{
magma_sdisplace_pointers(dW2_displ, T_array, ldt, i, j, batchCount, queue);
magmablas_slarft_recstrmv_sm32x32_batched(rows, mycol, dW3_displ, dW2_displ, ldt, dW1_displ, ldtstep, batchCount, queue);
}
}
}
// the upper rectangular protion is updated, now if needed update the triangular portion
if (stair_T == 0) {
if (DEBUG == 3) {
printf("doing strmv on the triangular portion of size %d %d of T(%d,%d)\n",
(int) mycol, (int) mycol, (int) j, (int) j );
}
if (mycol > 0)
{
magma_sdisplace_pointers(dW1_displ, dTstep_array, ldtstep, j, j, batchCount, queue);
magma_sdisplace_pointers(dW3_displ, tau_array, 1, j, 0, batchCount, queue);
magma_sdisplace_pointers(dW2_displ, T_array, ldt, j, j, batchCount, queue);
magmablas_slarft_strmv_sm32x32_batched(mycol, mycol, dW3_displ, dW1_displ, ldtstep, dW2_displ, ldt, batchCount, queue);
}
}
}// end of j
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dTstep_array);
return 0;
}
