extern "C" magma_int_t
magma_sgetrf_panel_nopiv_batched_q(
magma_int_t m, magma_int_t nb,
float** dA_array, magma_int_t ldda,
float** dX_array, magma_int_t dX_length,
float** dinvA_array, magma_int_t dinvA_length,
float** dW0_displ, float** dW1_displ,
float** dW2_displ, float** dW3_displ,
float** dW4_displ,
magma_int_t *info_array, magma_int_t gbstep,
magma_int_t batchCount, magma_queue_t stream, cublasHandle_t myhandle)
{
magma_int_t arginfo = 0;
//===============================================
// panel factorization
//===============================================
if(m < nb){
printf("magma_sgetrf_panel_nopiv_batched_q m < nb %d < %d \n",(int) m, (int) nb);
return -101;
}
#if 0
arginfo = magma_sgetf2_nopiv_batched(
m, nb,
dA_array, ldda,
dW1_displ, dW2_displ, dW3_displ,
info_array, gbstep, batchCount, myhandle);
if (arginfo != 0) return arginfo;
#else
arginfo = magma_sgetf2_nopiv_batched(
nb, nb,
dA_array, ldda,
dW1_displ, dW2_displ, dW3_displ,
info_array, gbstep, batchCount, myhandle);
if (arginfo != 0) return arginfo;
if((m-nb) > 0){
magma_sdisplace_pointers(dW0_displ, dA_array, ldda, nb, 0, batchCount);
magmablas_strsm_work_batched(MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
1, m-nb, nb,
MAGMA_S_ONE,
dA_array, ldda,
dW0_displ, ldda,
dX_array, m-nb,
dinvA_array, dinvA_length,
dW1_displ, dW2_displ,
dW3_displ, dW4_displ,
1, batchCount);
}
#endif
return 0;
}
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;
}
extern "C" magma_int_t
magma_sgetrf_recpanel_batched_q(
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 stream, cublasHandle_t myhandle)
{
//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, myhandle);
}
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);
magma_idisplace_pointers(dipiv_displ, dipiv_array, 1, p1, 0, batchCount);
magma_sgetrf_recpanel_batched_q(
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, stream, myhandle);
// update A2
//if(DEBUG>0)printf("calling TRSM with m=%d n=%d \n",m1,n2);
// setup pivinfo
setup_pivinfo_batched_q(dpivinfo_array, dipiv_displ, m1, n1, stream, batchCount);
magma_sdisplace_pointers(dW5_displ, dA_array, ldda, p1, p2, batchCount);
magma_slaswp_rowparallel_batched( n2, dW5_displ, ldda,
dX_array, n1,
0, n1,
dpivinfo_array, batchCount);
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);
magma_sdisplace_pointers(dW1_displ, dA_array, ldda, p2, 0, batchCount);
magma_sdisplace_pointers(dA_displ, dA_array, ldda, p2, p2, batchCount);
//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);
#if 0
float neg_one = MAGMA_S_NEG_ONE;
float one = MAGMA_S_ONE;
cublasSgemmBatched(myhandle, CUBLAS_OP_N, CUBLAS_OP_N, m2, n2, n1,
&neg_one, (const float**) dW1_displ, ldda,
(const float**) dW5_displ, ldda,
&one, dA_displ, ldda, batchCount );
#else
magmablas_sgemm_batched( MagmaNoTrans, MagmaNoTrans, m2, n2, n1,
MAGMA_S_NEG_ONE, dW1_displ, ldda,
dW5_displ, ldda,
MAGMA_S_ONE, dA_displ, ldda,
batchCount);
#endif
// 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);
magma_sgetrf_recpanel_batched_q(
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, stream, myhandle);
// setup pivinfo
setup_pivinfo_batched_q(dpivinfo_array, dipiv_displ, m2, n2, stream, batchCount);
adjust_ipiv_batched_q(dipiv_displ, n2, n1, magma_stream, batchCount);
magma_sdisplace_pointers(dW1_displ, dA_array, ldda, p2, 0, batchCount); // no need since it is above
magma_slaswp_rowparallel_batched( n1, dW1_displ, ldda,
dW1_displ, ldda,
n1, n,
dpivinfo_array, batchCount);
}
magma_free(dA_displ);
magma_free(dipiv_displ);
return 0;
}
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, cublasHandle_t myhandle,
magma_queue_t queue)
{
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
#ifdef RFT_MAG_GEM
magmablas_sgemm_batched( MagmaConjTrans, MagmaNoTrans,
k, k, n,
one, v_array, ldv,
v_array, ldv,
zero, dTstep_array, ldtstep,
batchCount, queue);
#else
cublasSgemmBatched(myhandle, CUBLAS_OP_C, CUBLAS_OP_N, k, k, n,
&one, (const float**) v_array, ldv,
(const float**) v_array, ldv,
&zero, dTstep_array, ldtstep, batchCount);
#endif
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);
//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",prev_n,mycol,0,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);
#ifdef RFT_MAG_GEM
magmablas_sgemm_batched( MagmaNoTrans, MagmaNoTrans,
prev_n, mycol, prev_n,
one, T_array, ldt,
dW1_displ, ldtstep,
zero, dW2_displ, ldt,
batchCount, queue );
#else
cublasSgemmBatched(myhandle, CUBLAS_OP_N, CUBLAS_OP_N,
prev_n, mycol, prev_n,
&one, (const float**) T_array, ldt,
(const float**) dW1_displ, ldtstep,
&zero, dW2_displ, ldt, batchCount);
#endif
// 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",rows,mycol,i,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",mycol,mycol,j,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;
}
extern "C" magma_int_t
magma_sgetrf_recpanel_nopiv_batched_q(
magma_int_t m, magma_int_t n, magma_int_t min_recpnb,
float** dA_array, magma_int_t ldda,
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 stream, cublasHandle_t myhandle)
{
// Quick return if possible
if (m == 0 || n == 0) {
return 0;
}
magma_int_t arginfo = 0;
float **dA_displ = NULL;
magma_malloc((void**)&dA_displ, batchCount * sizeof(*dA_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);
arginfo = magma_sgetrf_panel_nopiv_batched_q(
m, panel_nb,
dA_array, ldda,
dX_array, dX_length,
dinvA_array, dinvA_length,
dW1_displ, dW2_displ,
dW3_displ, dW4_displ, dW5_displ,
info_array, gbstep, batchCount, stream, myhandle);
if (arginfo != 0) return arginfo;
}
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
//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);
arginfo = magma_sgetrf_recpanel_nopiv_batched_q(
m1, n1, min_recpnb,
dA_displ, ldda,
dX_array, dX_length,
dinvA_array, dinvA_length,
dW1_displ, dW2_displ,
dW3_displ, dW4_displ, dW5_displ,
info_array, gbstep, batchCount, stream, myhandle);
if (arginfo != 0) return arginfo;
// update A2
//printf("calling update A2 with m=%d n=%d k=%d\n",m2,n2,n1);
magma_sdisplace_pointers(dW5_displ, dA_array, ldda, p1, p2, batchCount);
magmablas_strsm_work_batched(MagmaLeft, MagmaLower, MagmaNoTrans, MagmaUnit, 1,
n1, n2,
MAGMA_S_ONE,
dA_displ, ldda, // dA
dW5_displ, ldda, // dB
dX_array, n1, // dX
dinvA_array, dinvA_length,
dW1_displ, dW2_displ,
dW3_displ, dW4_displ,
1, batchCount);
magma_sdisplace_pointers(dW1_displ, dA_array, ldda, p2, 0, batchCount);
magma_sdisplace_pointers(dA_displ, dA_array, ldda, p2, p2, batchCount);
magmablas_sgemm_batched( MagmaNoTrans, MagmaNoTrans, m2, n2, n1,
MAGMA_S_NEG_ONE, dW1_displ, ldda,
dW5_displ, ldda,
MAGMA_S_ONE, dA_displ, ldda,
batchCount);
// panel on A2
//printf("calling recursive panel on A2 with m=%d nb=%d min_recpnb %d\n",m2,n2,min_recpnb);
arginfo = magma_sgetrf_recpanel_nopiv_batched_q(
m2, n2, min_recpnb,
dA_displ, ldda,
dX_array, dX_length,
dinvA_array, dinvA_length,
dW1_displ, dW2_displ,
dW3_displ, dW4_displ, dW5_displ,
info_array, gbstep+p2, batchCount, stream, myhandle);
if (arginfo != 0) return arginfo;
}
magma_free(dA_displ);
return 0;
}
extern "C" magma_int_t
magma_sgeqrf_expert_batched(
magma_int_t m, magma_int_t n,
float **dA_array, magma_int_t ldda,
float **dR_array, magma_int_t lddr,
float **dT_array, magma_int_t lddt,
float **dtau_array, magma_int_t provide_RT,
magma_int_t *info_array, magma_int_t batchCount, magma_queue_t queue)
{
#define dA(i, j) (dA + (i) + (j)*ldda) // A(i, j) means at i row, j column
/* Local Parameter */
magma_int_t nb = magma_get_sgeqrf_batched_nb(m);
magma_int_t nnb = 8;
magma_int_t min_mn = min(m, n);
/* Check arguments */
cudaMemset(info_array, 0, batchCount*sizeof(magma_int_t));
magma_int_t arginfo = 0;
if (m < 0)
arginfo = -1;
else if (n < 0)
arginfo = -2;
else if (ldda < max(1,m))
arginfo = -4;
else if (lddr < min_mn && provide_RT == 1)
arginfo = -6;
else if (lddr < min(min_mn, nb))
arginfo = -6;
else if (lddt < min(min_mn, nb))
arginfo = -8;
if (arginfo != 0) {
magma_xerbla( __func__, -(arginfo) );
return arginfo;
}
/* Quick return if possible */
if (m == 0 || n == 0)
if (min_mn == 0 ) return arginfo;
if ( m > 2048 || n > 2048 ) {
printf("=========================================================================================\n");
printf(" WARNING batched routines are designed for small sizes it might be better to use the\n Native/Hybrid classical routines if you want performance\n");
printf("=========================================================================================\n");
}
magma_int_t i, k, ib=nb, jb=nnb, offset_RT=0, use_stream;
magma_int_t ldw, offset;
float **dW0_displ = NULL;
float **dW1_displ = NULL;
float **dW2_displ = NULL;
float **dW3_displ = NULL;
float **dW4_displ = NULL;
float **dW5_displ = NULL;
float **dR_displ = NULL;
float **dT_displ = NULL;
float *dwork = NULL;
float **cpuAarray = NULL;
float **cpuTarray = NULL;
magma_malloc((void**)&dW0_displ, batchCount * sizeof(*dW0_displ));
magma_malloc((void**)&dW1_displ, batchCount * sizeof(*dW1_displ));
magma_malloc((void**)&dW2_displ, batchCount * sizeof(*dW2_displ));
magma_malloc((void**)&dW3_displ, batchCount * sizeof(*dW3_displ));
magma_malloc((void**)&dW4_displ, batchCount * sizeof(*dW4_displ));
magma_malloc((void**)&dW5_displ, batchCount * sizeof(*dW5_displ));
magma_malloc((void**)&dR_displ, batchCount * sizeof(*dR_displ));
magma_malloc((void**)&dT_displ, batchCount * sizeof(*dT_displ));
magma_smalloc(&dwork, (2 * nb * n) * batchCount);
magma_malloc_cpu((void**) &cpuAarray, batchCount*sizeof(float*));
magma_malloc_cpu((void**) &cpuTarray, batchCount*sizeof(float*));
/* check allocation */
if ( dW0_displ == NULL || dW1_displ == NULL || dW2_displ == NULL ||
dW3_displ == NULL || dW4_displ == NULL || dW5_displ == NULL ||
dR_displ == NULL || dT_displ == NULL || dwork == NULL ||
cpuAarray == NULL || cpuTarray == NULL ) {
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dW5_displ);
magma_free(dR_displ);
magma_free(dT_displ);
magma_free(dwork);
magma_free_cpu(cpuAarray);
magma_free_cpu(cpuTarray);
magma_int_t info = MAGMA_ERR_DEVICE_ALLOC;
magma_xerbla( __func__, -(info) );
return info;
}
magma_sdisplace_pointers(dR_displ, dR_array, lddr, 0, 0, batchCount, queue);
magma_sdisplace_pointers(dT_displ, dT_array, lddt, 0, 0, batchCount, queue);
// set dR and dT to zero. if provide_RT == 0 only a tile of size nbxnb is used and overwritten at each step
magmablas_slaset_batched( MagmaFull, lddr, (provide_RT > 0 ? n:min(min_mn,nb)), MAGMA_S_ZERO, MAGMA_S_ZERO, dR_displ, lddr, batchCount, queue );
magmablas_slaset_batched( MagmaFull, lddt, (provide_RT > 0 ? n:min(min_mn,nb)), MAGMA_S_ZERO, MAGMA_S_ZERO, dT_displ, lddt, batchCount, queue );
/*
if ( provide_RT > 0 )
{
magmablas_slaset_q( MagmaFull, lddr, n*batchCount, MAGMA_S_ZERO, MAGMA_S_ZERO, dR, lddr, queue );
magmablas_slaset_q( MagmaFull, lddt, n*batchCount, MAGMA_S_ZERO, MAGMA_S_ZERO, dT, lddt, queue );
}
else
{
magmablas_slaset_q( MagmaFull, lddr, nb*batchCount, MAGMA_S_ZERO, MAGMA_S_ZERO, dR, lddr, queue );
magmablas_slaset_q( MagmaFull, lddt, nb*batchCount, MAGMA_S_ZERO, MAGMA_S_ZERO, dT, lddt, queue );
}
*/
magma_int_t streamid;
const magma_int_t nbstreams=10;
magma_queue_t queues[nbstreams];
for (i=0; i < nbstreams; i++) {
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[i] );
}
magma_getvector( batchCount, sizeof(float*), dA_array, 1, cpuAarray, 1, queue);
magma_getvector( batchCount, sizeof(float*), dT_array, 1, cpuTarray, 1, queue);
for (i=0; i < min_mn; i += nb)
{
ib = min(nb, min_mn-i);
//===============================================
// panel factorization
//===============================================
magma_sdisplace_pointers(dW0_displ, dA_array, ldda, i, i, batchCount, queue);
magma_sdisplace_pointers(dW2_displ, dtau_array, 1, i, 0, batchCount, queue);
if ( provide_RT > 0 )
{
offset_RT = i;
magma_sdisplace_pointers(dR_displ, dR_array, lddr, (provide_RT == 1 ? offset_RT:0), offset_RT, batchCount, queue);
magma_sdisplace_pointers(dT_displ, dT_array, lddt, 0, offset_RT, batchCount, queue);
}
//dwork is used in panel factorization and trailing matrix update
//dW4_displ, dW5_displ are used as workspace and configured inside
magma_sgeqrf_panel_batched(m-i, ib, jb,
dW0_displ, ldda,
dW2_displ,
dT_displ, lddt,
dR_displ, lddr,
dW1_displ,
dW3_displ,
dwork,
dW4_displ, dW5_displ,
info_array,
batchCount, queue);
//===============================================
// end of panel
//===============================================
//===============================================
// update trailing matrix
//===============================================
if ( (n-ib-i) > 0)
{
//dwork is used in panel factorization and trailing matrix update
//reset dW4_displ
ldw = nb;
magma_sset_pointer( dW4_displ, dwork, 1, 0, 0, ldw*n, batchCount, queue );
offset = ldw*n*batchCount;
magma_sset_pointer( dW5_displ, dwork + offset, 1, 0, 0, ldw*n, batchCount, queue );
// set the diagonal of v as one and the upper triangular part as zero already set inside geqrf_panel
//magmablas_slaset_batched( MagmaUpper, ib, ib, MAGMA_S_ZERO, MAGMA_S_ONE, dW0_displ, ldda, batchCount, queue );
//magma_sdisplace_pointers(dW2_displ, dtau_array, 1, i, 0, batchCount, queue);
// it is faster since it is using BLAS-3 GEMM routines, different from lapack implementation
magma_slarft_batched(m-i, ib, 0,
dW0_displ, ldda,
dW2_displ,
dT_displ, lddt,
dW4_displ, nb*lddt,
batchCount, queue);
// perform C = (I-V T^H V^H) * C, C is the trailing matrix
//-------------------------------------------
// USE STREAM GEMM
//-------------------------------------------
use_stream = magma_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;
// the queue gemm must take cpu pointer
magma_slarfb_gpu_gemm( MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-i, n-i-ib, ib,
cpuAarray[k] + i + i * ldda, ldda,
cpuTarray[k] + offset_RT*lddt, lddt,
cpuAarray[k] + i + (i+ib) * ldda, ldda,
dwork + nb * n * k, -1,
dwork + nb * n * batchCount + nb * n * k, -1, queues[streamid] );
}
// need to synchronise to be sure that panel does not start before
// finishing the update at least of the next panel
// if queue is NULL, no need to sync
if ( queue != NULL ) {
for (magma_int_t s=0; s < nbstreams; s++)
magma_queue_sync(queues[s]);
}
}
//-------------------------------------------
// USE BATCHED GEMM
//-------------------------------------------
else
{
//direct trailing matrix in dW1_displ
magma_sdisplace_pointers(dW1_displ, dA_array, ldda, i, i+ib, batchCount, queue);
magma_slarfb_gemm_batched(
MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-i, n-i-ib, ib,
(const float**)dW0_displ, ldda,
(const float**)dT_displ, lddt,
dW1_displ, ldda,
dW4_displ, ldw,
dW5_displ, ldw,
batchCount, queue );
}
}// update the trailing matrix
//===============================================
// copy dR back to V after the trailing matrix update,
// only when provide_RT=0 otherwise the nbxnb block of V is set to diag=1/0
// The upper portion of V could be set totaly to 0 here
if ( provide_RT == 0 )
{
magmablas_slacpy_batched( MagmaUpper, ib, ib, dR_displ, lddr, dW0_displ, ldda, batchCount, queue );
}
}
magma_queue_sync(queue);
for (k=0; k < nbstreams; k++) {
magma_queue_destroy( queues[k] );
}
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dW5_displ);
magma_free(dR_displ);
magma_free(dT_displ);
magma_free(dwork);
magma_free_cpu(cpuAarray);
magma_free_cpu(cpuTarray);
return arginfo;
}
/**
Purpose
-------
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;
}
extern "C" magma_int_t
magma_sgetf2_batched(
magma_int_t m, magma_int_t n,
float **dA_array, magma_int_t lda,
float **dW0_displ,
float **dW1_displ,
float **dW2_displ,
magma_int_t **ipiv_array,
magma_int_t *info_array,
magma_int_t gbstep,
magma_int_t batchCount,
cublasHandle_t myhandle, magma_queue_t queue)
{
magma_int_t arginfo = 0;
if (m < 0) {
arginfo = -1;
} else if (n < 0 ) {
arginfo = -2;
} else if (lda < 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 neg_one = MAGMA_S_NEG_ONE;
float one = MAGMA_S_ONE;
magma_int_t nb = BATF2_NB;
//float **cpuAarray = (float**) malloc(batchCount*sizeof(float*));
//magma_getvector( batchCount, sizeof(float*), dA_array, 1, cpuAarray, 1);
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;
//size_t required_shmem_size = zamax*(sizeof(float)+sizeof(int)) + (m-panelj+2)*sizeof(float);
//if( (m-panelj) > 0)
if( (m-panelj) > MAX_NTHREADS)
//if( required_shmem_size > (MAX_SHARED_ALLOWED*1024))
{
//printf("running non shared version\n");
// find the max of the column gbj
arginfo = magma_isamax_batched(m-gbj, dA_array, 1, gbj, lda, ipiv_array, info_array, gbstep, batchCount, queue);
if(arginfo != 0 ) return arginfo;
// Apply the interchange to columns 1:N. swap the whole row
arginfo = magma_sswap_batched(n, dA_array, lda, gbj, ipiv_array, batchCount, queue);
if(arginfo != 0 ) return arginfo;
// Compute elements J+1:M of J-th column.
if (gbj < m) {
arginfo = magma_sscal_sger_batched(m-gbj, ib-step, gbj, dA_array, lda, info_array, gbstep, batchCount, queue);
if(arginfo != 0 ) return arginfo;
}
}
else{
//printf("running --- shared version\n");
arginfo = magma_scomputecolumn_batched(m-panelj, panelj, step, dA_array, lda, ipiv_array, info_array, gbstep, batchCount, queue);
if(arginfo != 0 ) return arginfo;
// Apply the interchange to columns 1:N. swap the whole row
arginfo = magma_sswap_batched(n, dA_array, lda, gbj, ipiv_array, batchCount, queue);
if(arginfo != 0 ) return arginfo;
}
}
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, lda, batchCount, queue);
// do the blocked DGER = DGEMM for the remaining panelj+ib:n columns
magma_sdisplace_pointers(dW0_displ, dA_array, lda, ib+panelj, panelj, batchCount, queue);
magma_sdisplace_pointers(dW1_displ, dA_array, lda, panelj, ib+panelj, batchCount, queue);
magma_sdisplace_pointers(dW2_displ, dA_array, lda, ib+panelj, ib+panelj, batchCount, queue);
#if 1
magmablas_sgemm_batched( MagmaNoTrans, MagmaNoTrans, m-(panelj+ib), n-(panelj+ib), ib,
neg_one, dW0_displ, lda,
dW1_displ, lda,
one, dW2_displ, lda,
batchCount, queue);
#else
cublasSgemmBatched(myhandle, CUBLAS_OP_N, CUBLAS_OP_N, m-(panelj+ib), n-(panelj+ib), ib,
&neg_one, (const float**) dW0_displ, lda,
(const float**) dW1_displ, lda,
&one, dW2_displ, lda, batchCount );
#endif
}
}
//free(cpuAarray);
return 0;
}
extern "C" magma_int_t
magma_sgeqrf_panel_batched(
magma_int_t m, magma_int_t n, magma_int_t nb,
float** dA_array, magma_int_t ldda,
float** tau_array,
float** dT_array, magma_int_t ldt,
float** dR_array, magma_int_t ldr,
float** dW0_displ,
float** dW1_displ,
float *dwork,
float** dW2_displ,
float** 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_sdisplace_pointers(dW0_displ, dA_array, ldda, j, j, batchCount, queue);
magma_sdisplace_pointers(dW2_displ, tau_array, 1, j, 0, batchCount, queue);
magma_sdisplace_pointers(dW3_displ, dR_array, ldr, j, j, batchCount, queue); //
//sub-panel factorization
magma_sgeqr2_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_sdisplace_pointers(dW0_displ, dA_array, ldda, 0, j, batchCount, queue);
magma_sdisplace_pointers(dW3_displ, dR_array, ldr, 0, j, batchCount, queue);
magmablas_slacpy_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_sdisplace_pointers(dW0_displ, dA_array, ldda, j, j, batchCount, queue);
magmablas_slaset_batched( MagmaUpper, jb, jb, MAGMA_S_ZERO, MAGMA_S_ONE, dW0_displ, ldda, batchCount, queue );
if ( (n-j-jb) > 0) //update the trailing matrix inside the panel
{
magma_slarft_sm32x32_batched(m-j, jb,
dW0_displ, ldda,
dW2_displ,
dT_array, ldt,
batchCount, queue);
magma_sdisplace_pointers( dW1_displ, dA_array, ldda, j, j + jb, batchCount, queue );
magma_sset_pointer( dW2_displ, dwork, 1, 0, 0, ldw*n, batchCount, queue );
magma_sset_pointer( dW3_displ, dwork + ldw*n*batchCount, 1, 0, 0, ldw*n, batchCount, queue );
magma_slarfb_gemm_batched(
MagmaLeft, MagmaConjTrans, MagmaForward, MagmaColumnwise,
m-j, n-j-jb, jb,
(const float**)dW0_displ, ldda,
(const float**)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_sdisplace_pointers(dW0_displ, dA_array, ldda, 0, minmn, batchCount, queue);
magma_sdisplace_pointers(dW3_displ, dR_array, ldr, 0, minmn, batchCount, queue);
magmablas_slacpy_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 sgeqrf_batched
magma_sdisplace_pointers(dW0_displ, dA_array, ldda, 0, 0, batchCount, queue);
magmablas_slaset_batched( MagmaUpper, minmn, n, MAGMA_S_ZERO, MAGMA_S_ONE, dW0_displ, ldda, batchCount, queue );
return MAGMA_SUCCESS;
}
/**
Purpose
-------
SPOTRF computes the Cholesky factorization of a real symmetric
positive definite matrix dA.
The factorization has the form
dA = U**H * U, if UPLO = MagmaUpper, or
dA = L * L**H, if UPLO = MagmaLower,
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
If the current stream is NULL, this version replaces it with a new
stream to overlap computation with communication.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of dA is stored;
- = MagmaLower: Lower triangle of dA is stored.
@param[in]
n INTEGER
The order of the matrix dA. N >= 0.
@param[in,out]
dA REAL array on the GPU, dimension (LDDA,N)
On entry, the symmetric matrix dA. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of dA contains the upper
triangular part of the matrix dA, and the strictly lower
triangular part of dA is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of dA contains the lower
triangular part of the matrix dA, and the strictly upper
triangular part of dA is not referenced.
\n
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization dA = U**H * U or dA = L * L**H.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
@ingroup magma_sposv_comp
********************************************************************/
extern "C" magma_int_t
magma_spotrf_batched(
magma_uplo_t uplo, magma_int_t n,
float **dA_array, magma_int_t ldda,
magma_int_t *info_array, magma_int_t batchCount)
{
#define A(i_, j_) (A + (i_) + (j_)*ldda)
cudaMemset(info_array, 0, batchCount*sizeof(magma_int_t));
magma_int_t arginfo = 0;
if ( uplo != MagmaUpper && uplo != MagmaLower) {
arginfo = -1;
} else if (n < 0) {
arginfo = -2;
} else if (ldda < max(1,n)) {
arginfo = -4;
}
if (arginfo != 0) {
magma_xerbla( __func__, -(arginfo) );
return arginfo;
}
// Quick return if possible
if (n == 0) {
return arginfo;
}
if( n > 2048 ){
printf("=========================================================================================\n");
printf(" WARNING batched routines are designed for small sizes it might be better to use the\n Native/Hybrid classical routines if you want performance\n");
printf("=========================================================================================\n");
}
magma_int_t j, k, ib;
magma_int_t nb = POTRF_NB;
magma_int_t gemm_crossover = 127;//nb > 32 ? 127 : 160;
#if defined(USE_CUOPT)
cublasHandle_t myhandle;
cublasCreate_v2(&myhandle);
#else
cublasHandle_t myhandle=NULL;
#endif
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 **dx_array = NULL;
magma_malloc((void**)&dA_displ, batchCount * sizeof(*dA_displ));
magma_malloc((void**)&dW0_displ, batchCount * sizeof(*dW0_displ));
magma_malloc((void**)&dW1_displ, batchCount * sizeof(*dW1_displ));
magma_malloc((void**)&dW2_displ, batchCount * sizeof(*dW2_displ));
magma_malloc((void**)&dW3_displ, batchCount * sizeof(*dW3_displ));
magma_malloc((void**)&dW4_displ, batchCount * sizeof(*dW4_displ));
magma_malloc((void**)&dinvA_array, batchCount * sizeof(*dinvA_array));
magma_malloc((void**)&dx_array, batchCount * sizeof(*dx_array));
float* dinvA;
float* dx;// dinvA and x are workspace in strsm
magma_int_t invA_msize = ((n+TRI_NB-1)/TRI_NB)*TRI_NB*TRI_NB;
magma_int_t x_msize = n*nb;
magma_smalloc( &dinvA, invA_msize * batchCount);
magma_smalloc( &dx, x_msize * batchCount );
sset_pointer(dx_array, dx, 1, 0, 0, x_msize, batchCount);
sset_pointer(dinvA_array, dinvA, TRI_NB, 0, 0, invA_msize, batchCount);
cudaMemset( dinvA, 0, batchCount * ((n+TRI_NB-1)/TRI_NB)*TRI_NB*TRI_NB * sizeof(float) );
float **cpuAarray = NULL;
magma_malloc_cpu((void**) &cpuAarray, batchCount*sizeof(float*));
magma_getvector( batchCount, sizeof(float*), dA_array, 1, cpuAarray, 1);
float d_alpha = -1.0;
float d_beta = 1.0;
magma_queue_t cstream;
magmablasGetKernelStream(&cstream);
magma_int_t streamid;
const magma_int_t nbstreams=32;
magma_queue_t stream[nbstreams];
for(k=0; k<nbstreams; k++){
magma_queue_create( &stream[k] );
}
magmablasSetKernelStream(NULL);
if (uplo == MagmaUpper) {
printf("Upper side is unavailable \n");
goto fin;
}
else {
for(j = 0; j < n; j+=nb) {
ib = min(nb, n-j);
#if 1
//===============================================
// panel factorization
//===============================================
magma_sdisplace_pointers(dA_displ, dA_array, ldda, j, j, batchCount);
sset_pointer(dx_array, dx, 1, 0, 0, x_msize, batchCount);
sset_pointer(dinvA_array, dinvA, TRI_NB, 0, 0, invA_msize, batchCount);
#if 0
arginfo = magma_spotrf_panel_batched(
uplo, n-j, ib,
dA_displ, ldda,
dx_array, x_msize,
dinvA_array, invA_msize,
dW0_displ, dW1_displ, dW2_displ,
dW3_displ, dW4_displ,
info_array, j, batchCount, myhandle);
#else
//arginfo = magma_spotrf_rectile_batched(
arginfo = magma_spotrf_recpanel_batched(
uplo, n-j, ib, 32,
dA_displ, ldda,
dx_array, x_msize,
dinvA_array, invA_msize,
dW0_displ, dW1_displ, dW2_displ,
dW3_displ, dW4_displ,
info_array, j, batchCount, myhandle);
#endif
if(arginfo != 0 ) goto fin;
//===============================================
// end of panel
//===============================================
#endif
#if 1
//real_Double_t gpu_time;
//gpu_time = magma_sync_wtime(NULL);
if( (n-j-ib) > 0){
if( (n-j-ib) > gemm_crossover)
{
//-------------------------------------------
// USE STREAM HERK
//-------------------------------------------
// since it use different stream I need to wait the panel.
// But since the code use the NULL stream everywhere,
// so I don't need it, because the NULL stream do the sync by itself
//magma_queue_sync(NULL);
/* you must know the matrix layout inorder to do it */
for(k=0; k<batchCount; k++)
{
streamid = k%nbstreams;
magmablasSetKernelStream(stream[streamid]);
// call herk, class ssyrk must call cpu pointer
magma_ssyrk(MagmaLower, MagmaNoTrans, n-j-ib, ib,
d_alpha,
(const float*) cpuAarray[k] + j+ib+j*ldda, ldda,
d_beta,
cpuAarray[k] + j+ib+(j+ib)*ldda, ldda);
}
// need to synchronise to be sure that panel do not start before
// finishing the update at least of the next panel
// BUT no need for it as soon as the other portion of the code
// use the NULL stream which do the sync by itself
//magma_device_sync();
magmablasSetKernelStream(NULL);
}
else
{
//-------------------------------------------
// USE BATCHED GEMM(which is a HERK in fact, since it only access the lower part)
//-------------------------------------------
magma_sdisplace_pointers(dA_displ, dA_array, ldda, j+ib, j, batchCount);
magma_sdisplace_pointers(dW1_displ, dA_array, ldda, j+ib, j+ib, batchCount);
magmablas_ssyrk_batched(uplo, MagmaNoTrans, n-j-ib, ib,
d_alpha, dA_displ, ldda,
d_beta, dW1_displ, ldda,
batchCount);
}
}
//gpu_time = magma_sync_wtime(NULL) - gpu_time;
//real_Double_t flops = (n-j-ib) * (n-j-ib) * ib / 1e9 * batchCount;
//real_Double_t gpu_perf = flops / gpu_time;
//printf("Rows= %d, Colum=%d, herk time = %7.2fms, Gflops= %7.2f\n", n-j-ib, ib, gpu_time*1000, gpu_perf);
#endif
}
}
fin:
magma_queue_sync(NULL);
for(k=0; k<nbstreams; k++){
magma_queue_destroy( stream[k] );
}
magmablasSetKernelStream(cstream);
#if defined(USE_CUOPT)
cublasDestroy_v2(myhandle);
#endif
magma_free(dA_displ);
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dinvA_array);
magma_free(dx_array);
magma_free(dinvA);
magma_free(dx);
magma_free_cpu(cpuAarray);
return arginfo;
}
/***************************************************************************//**
Purpose
-------
SPOTRF computes the Cholesky factorization of a real symmetric
positive definite matrix dA.
The factorization has the form
dA = U**H * U, if UPLO = MagmaUpper, or
dA = L * L**H, if UPLO = MagmaLower,
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of dA is stored;
- = MagmaLower: Lower triangle of dA is stored.
@param[in]
n INTEGER
The order of the matrix dA. N >= 0.
@param[in,out]
dA_array Array of pointers, dimension (batchCount).
Each is a REAL array on the GPU, dimension (LDDA,N)
On entry, each pointer is a symmetric matrix dA.
If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of dA contains the upper
triangular part of the matrix dA, and the strictly lower
triangular part of dA is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of dA contains the lower
triangular part of the matrix dA, and the strictly upper
triangular part of dA is not referenced.
\n
On exit, if corresponding entry in info_array = 0,
each pointer is the factor U or L from the Cholesky
factorization dA = U**H * U or dA = L * L**H.
@param[in]
ldda INTEGER
The leading dimension of each array dA. LDDA >= max(1,N).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param[out]
info_array Array of INTEGERs, dimension (batchCount), for corresponding matrices.
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
@param[in]
batchCount INTEGER
The number of matrices to operate on.
@param[in]
queue magma_queue_t
Queue to execute in.
@ingroup magma_potrf_batched
*******************************************************************************/
extern "C" magma_int_t
magma_spotrf_lg_batched(
magma_uplo_t uplo, magma_int_t n,
float **dA_array, magma_int_t ldda,
magma_int_t *info_array, magma_int_t batchCount, magma_queue_t queue)
{
magma_int_t arginfo = 0;
#define A(i_, j_) (A + (i_) + (j_)*ldda)
float d_alpha = -1.0;
float d_beta = 1.0;
if ( n > 2048 ) {
#ifndef MAGMA_NOWARNING
printf("=========================================================================================\n"
" WARNING batched routines are designed for small sizes. It might be better to use the\n"
" Native/Hybrid classical routines if you want good performance.\n"
"=========================================================================================\n");
#endif
}
magma_int_t j, k, ib, use_stream;
magma_int_t nb, recnb;
magma_get_spotrf_batched_nbparam(n, &nb, &recnb);
float **dA_displ = NULL;
float **dW0_displ = NULL;
float **dW1_displ = NULL;
float **dW2_displ = NULL;
float **dW3_displ = NULL;
float **dW4_displ = NULL;
float **dinvA_array = NULL;
float **dwork_array = NULL;
magma_malloc((void**)&dA_displ, batchCount * sizeof(*dA_displ));
magma_malloc((void**)&dW0_displ, batchCount * sizeof(*dW0_displ));
magma_malloc((void**)&dW1_displ, batchCount * sizeof(*dW1_displ));
magma_malloc((void**)&dW2_displ, batchCount * sizeof(*dW2_displ));
magma_malloc((void**)&dW3_displ, batchCount * sizeof(*dW3_displ));
magma_malloc((void**)&dW4_displ, batchCount * sizeof(*dW4_displ));
magma_malloc((void**)&dinvA_array, batchCount * sizeof(*dinvA_array));
magma_malloc((void**)&dwork_array, batchCount * sizeof(*dwork_array));
magma_int_t invA_msize = magma_roundup( n, STRTRI_BATCHED_NB )*STRTRI_BATCHED_NB;
magma_int_t dwork_msize = n*nb;
float* dinvA = NULL;
float* dwork = NULL; // dinvA and dwork are workspace in strsm
float **cpuAarray = NULL;
magma_smalloc( &dinvA, invA_msize * batchCount);
magma_smalloc( &dwork, dwork_msize * batchCount );
magma_malloc_cpu((void**) &cpuAarray, batchCount*sizeof(float*));
/* check allocation */
if ( dA_displ == NULL || dW0_displ == NULL || dW1_displ == NULL || dW2_displ == NULL ||
dW3_displ == NULL || dW4_displ == NULL || dinvA_array == NULL || dwork_array == NULL ||
dinvA == NULL || dwork == NULL || cpuAarray == NULL ) {
magma_free(dA_displ);
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dinvA_array);
magma_free(dwork_array);
magma_free( dinvA );
magma_free( dwork );
magma_free_cpu(cpuAarray);
magma_int_t info = MAGMA_ERR_DEVICE_ALLOC;
magma_xerbla( __func__, -(info) );
return info;
}
magmablas_slaset_q( MagmaFull, invA_msize, batchCount, MAGMA_S_ZERO, MAGMA_S_ZERO, dinvA, invA_msize, queue );
magmablas_slaset_q( MagmaFull, dwork_msize, batchCount, MAGMA_S_ZERO, MAGMA_S_ZERO, dwork, dwork_msize, queue );
magma_sset_pointer( dwork_array, dwork, 1, 0, 0, dwork_msize, batchCount, queue );
magma_sset_pointer( dinvA_array, dinvA, STRTRI_BATCHED_NB, 0, 0, invA_msize, batchCount, queue );
magma_int_t streamid;
const magma_int_t nbstreams=10;
magma_queue_t queues[nbstreams];
for (k=0; k < nbstreams; k++) {
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[k] );
}
magma_getvector( batchCount, sizeof(float*), dA_array, 1, cpuAarray, 1, queue);
if (uplo == MagmaUpper) {
printf("Upper side is unavailable\n");
goto fin;
}
else {
for (j = 0; j < n; j += nb) {
ib = min(nb, n-j);
#if 1
//===============================================
// panel factorization
//===============================================
magma_sdisplace_pointers(dA_displ, dA_array, ldda, j, j, batchCount, queue);
magma_sset_pointer( dwork_array, dwork, 1, 0, 0, dwork_msize, batchCount, queue );
magma_sset_pointer( dinvA_array, dinvA, STRTRI_BATCHED_NB, 0, 0, invA_msize, batchCount, queue );
if (recnb == nb)
{
arginfo = magma_spotrf_panel_batched(
uplo, n-j, ib,
dA_displ, ldda,
dwork_array, dwork_msize,
dinvA_array, invA_msize,
dW0_displ, dW1_displ, dW2_displ,
dW3_displ, dW4_displ,
info_array, j, batchCount, queue);
}
else {
//arginfo = magma_spotrf_rectile_batched(
arginfo = magma_spotrf_recpanel_batched(
uplo, n-j, ib, recnb,
dA_displ, ldda,
dwork_array, dwork_msize,
dinvA_array, invA_msize,
dW0_displ, dW1_displ, dW2_displ,
dW3_displ, dW4_displ,
info_array, j, batchCount, queue);
}
if (arginfo != 0 ) goto fin;
//===============================================
// end of panel
//===============================================
#endif
#if 1
//real_Double_t gpu_time;
//gpu_time = magma_sync_wtime(queue);
if ( (n-j-ib) > 0) {
use_stream = magma_srecommend_cublas_gemm_stream(MagmaNoTrans, MagmaConjTrans, n-j-ib, n-j-ib, ib);
if (use_stream)
{
//-------------------------------------------
// USE STREAM HERK
//-------------------------------------------
// since it use different queue I need to wait the panel.
/* you must know the matrix layout inorder to do it */
magma_queue_sync(queue);
for (k=0; k < batchCount; k++)
{
streamid = k%nbstreams;
// call herk, class ssyrk must call cpu pointer
magma_ssyrk( MagmaLower, MagmaNoTrans, n-j-ib, ib,
d_alpha,
(const float*) cpuAarray[k] + j+ib+j*ldda, ldda,
d_beta,
cpuAarray[k] + j+ib+(j+ib)*ldda, ldda, queues[streamid] );
}
// need to synchronise to be sure that panel do not start before
// finishing the update at least of the next panel
// if queue is NULL, no need to sync
if ( queue != NULL ) {
for (magma_int_t s=0; s < nbstreams; s++)
magma_queue_sync(queues[s]);
}
}
else
{
//-------------------------------------------
// USE BATCHED GEMM(which is a HERK in fact, since it only access the lower part)
//-------------------------------------------
magma_sdisplace_pointers(dA_displ, dA_array, ldda, j+ib, j, batchCount, queue);
magma_sdisplace_pointers(dW1_displ, dA_array, ldda, j+ib, j+ib, batchCount, queue);
magmablas_ssyrk_batched( uplo, MagmaNoTrans, n-j-ib, ib,
d_alpha, dA_displ, ldda,
d_beta, dW1_displ, ldda,
batchCount, queue );
}
}
//gpu_time = magma_sync_wtime(queue) - gpu_time;
//real_Double_t flops = (n-j-ib) * (n-j-ib) * ib / 1e9 * batchCount;
//real_Double_t gpu_perf = flops / gpu_time;
//printf("Rows= %lld, Colum=%lld, herk time = %7.2fms, Gflops= %7.2f\n",
// (long long)(n-j-ib), (long long) ib, gpu_time*1000, gpu_perf);
#endif
}
}
fin:
magma_queue_sync(queue);
for (k=0; k < nbstreams; k++) {
magma_queue_destroy( queues[k] );
}
magma_free(dA_displ);
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dinvA_array);
magma_free(dwork_array);
magma_free( dinvA );
magma_free( dwork );
magma_free_cpu(cpuAarray);
return arginfo;
}
