/**
Purpose
-------
ZHERK performs one of the Hermitian rank k operations
C := alpha*A*A**H + beta*C,
or
C := alpha*A**H*A + beta*C,
where alpha and beta are real scalars, C is an n by n Hermitian
matrix and A is an n by k matrix in the first case and a k by n
matrix in the second case.
Parameters
----------
@param[in]
uplo magma_uplo_t.
On entry, uplo specifies whether the upper or lower
triangular part of the array C is to be referenced as
follows:
uplo = MagmaUpper Only the upper triangular part of C
is to be referenced.
uplo = MagmaLower Only the lower triangular part of C
is to be referenced.
@param[in]
trans magma_trans_t.
On entry, trans specifies the operation to be performed as
follows:
trans = MagmaNoTrans C := alpha*A*A**H + beta*C.
trans = MagmaConjTrans C := alpha*A**H*A + beta*C.
@param[in]
n INTEGER.
On entry, specifies the order of the matrix C. N must be
at least zero.
@param[in]
k INTEGER.
On entry with trans = MagmaNoTrans, k specifies the number
of columns of the matrix A, and on entry with
trans = MagmaConjTrans, k specifies the number of rows of the
matrix A. K must be at least zero.
@param[in]
alpha DOUBLE PRECISION
On entry, ALPHA specifies the scalar alpha.
@param[in]
dA_array Array of pointers, dimension (batchCount).
Each is a COMPLEX_16 A array of DIMENSION ( ldda, ka ), where ka is
k when trans = MagmaNoTrans, and is n otherwise.
Before entry with trans = MagmaNoTrans, the leading m by k
part of the array A must contain the matrix A, otherwise
the leading k by m part of the array A must contain the
matrix A.
@param[in]
ldda INTEGER.
On entry, ldda specifies the first dimension of each array A as declared
in the calling (sub) program. When trans = MagmaNoTrans then
ldda must be at least max( 1, n ), otherwise ldda must be at
least max( 1, k ).
@param[in]
beta DOUBLE PRECISION.
On entry, BETA specifies the scalar beta. When BETA is
supplied as zero then C need not be set on input.
@param[in,out]
dC_array Array of pointers, dimension (batchCount).
Each is a COMPLEX_16 array C of DIMENSION ( lddc, n ).
Before entry with uplo = MagmaUpper, the leading n by n
upper triangular part of the array C must contain the upper
triangular part of the Hermitian matrix and the strictly
lower triangular part of C is not referenced. On exit, the
upper triangular part of the array C is overwritten by the
upper triangular part of the updated matrix.
Before entry with uplo = MagmaLower, the leading n by n
lower triangular part of the array C must contain the lower
triangular part of the Hermitian matrix and the strictly
upper triangular part of C is not referenced. On exit, the
lower triangular part of the array C is overwritten by the
lower triangular part of the updated matrix.
Note that the imaginary parts of the diagonal elements need
not be set, they are assumed to be zero, and on exit they
are set to zero.
@param[in]
lddc INTEGER.
On entry, lddc specifies the first dimension of each array C as declared
in the calling (sub) program. lddc must be at least
max( 1, m ).
@param[in]
batchCount INTEGER
The number of matrices to operate on.
@param[in]
queue magma_queue_t
Queue to execute in.
@ingroup magma_zblas3
********************************************************************/
extern "C" void
magma_zherk_batched(
magma_uplo_t uplo, magma_trans_t trans, magma_int_t n, magma_int_t k,
double alpha,
magmaDoubleComplex const * const * dA_array, magma_int_t ldda,
double beta,
magmaDoubleComplex **dC_array, magma_int_t lddc, magma_int_t batchCount, magma_queue_t queue )
{
magmablas_zherk_batched(
uplo, trans, n, k,
alpha, dA_array, ldda,
beta, dC_array, lddc,
batchCount, queue );
}
/**
Purpose
-------
ZPOTRF computes the Cholesky factorization of a complex Hermitian
positive definite matrix dA.
The factorization has the form
dA = U**H * U, if UPLO = MagmaUpper, or
dA = L * L**H, if UPLO = MagmaLower,
where U is an upper triangular matrix and L is lower triangular.
This is the block version of the algorithm, calling Level 3 BLAS.
If the current stream is NULL, this version replaces it with a new
stream to overlap computation with communication.
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of dA is stored;
- = MagmaLower: Lower triangle of dA is stored.
@param[in]
n INTEGER
The order of the matrix dA. N >= 0.
@param[in,out]
dA COMPLEX_16 array on the GPU, dimension (LDDA,N)
On entry, the Hermitian matrix dA. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of dA contains the upper
triangular part of the matrix dA, and the strictly lower
triangular part of dA is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of dA contains the lower
triangular part of the matrix dA, and the strictly upper
triangular part of dA is not referenced.
\n
On exit, if INFO = 0, the factor U or L from the Cholesky
factorization dA = U**H * U or dA = L * L**H.
@param[in]
ldda INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
To benefit from coalescent memory accesses LDDA must be
divisible by 16.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
- > 0: if INFO = i, the leading minor of order i is not
positive definite, and the factorization could not be
completed.
@ingroup magma_zposv_comp
********************************************************************/
extern "C" magma_int_t
magma_zpotrf_batched(
magma_uplo_t uplo, magma_int_t n,
magmaDoubleComplex **dA_array, magma_int_t ldda,
magma_int_t *info_array, magma_int_t batchCount, magma_queue_t queue)
{
#define A(i_, j_) (A + (i_) + (j_)*ldda)
double d_alpha = -1.0;
double d_beta = 1.0;
cudaMemset(info_array, 0, batchCount*sizeof(magma_int_t));
magma_int_t arginfo = 0;
if ( uplo != MagmaUpper && uplo != MagmaLower) {
arginfo = -1;
} else if (n < 0) {
arginfo = -2;
} else if (ldda < max(1,n)) {
arginfo = -4;
}
if (arginfo != 0) {
magma_xerbla( __func__, -(arginfo) );
return arginfo;
}
// Quick return if possible
if (n == 0) {
return arginfo;
}
if( n > 2048 ){
printf("=========================================================================================\n");
printf(" WARNING batched routines are designed for small sizes it might be better to use the\n Native/Hybrid classical routines if you want performance\n");
printf("=========================================================================================\n");
}
magma_int_t j, k, ib;
magma_int_t nb = POTRF_NB;
magma_int_t gemm_crossover = 127;//nb > 32 ? 127 : 160;
#if defined(USE_CUOPT)
cublasHandle_t myhandle;
cublasCreate_v2(&myhandle);
#else
cublasHandle_t myhandle=NULL;
#endif
magmaDoubleComplex **dA_displ = NULL;
magmaDoubleComplex **dW0_displ = NULL;
magmaDoubleComplex **dW1_displ = NULL;
magmaDoubleComplex **dW2_displ = NULL;
magmaDoubleComplex **dW3_displ = NULL;
magmaDoubleComplex **dW4_displ = NULL;
magmaDoubleComplex **dinvA_array = NULL;
magmaDoubleComplex **dwork_array = NULL;
magma_malloc((void**)&dA_displ, batchCount * sizeof(*dA_displ));
magma_malloc((void**)&dW0_displ, batchCount * sizeof(*dW0_displ));
magma_malloc((void**)&dW1_displ, batchCount * sizeof(*dW1_displ));
magma_malloc((void**)&dW2_displ, batchCount * sizeof(*dW2_displ));
magma_malloc((void**)&dW3_displ, batchCount * sizeof(*dW3_displ));
magma_malloc((void**)&dW4_displ, batchCount * sizeof(*dW4_displ));
magma_malloc((void**)&dinvA_array, batchCount * sizeof(*dinvA_array));
magma_malloc((void**)&dwork_array, batchCount * sizeof(*dwork_array));
magma_int_t invA_msize = ((n+TRI_NB-1)/TRI_NB)*TRI_NB*TRI_NB;
magma_int_t dwork_msize = n*nb;
magmaDoubleComplex* dinvA = NULL;
magmaDoubleComplex* dwork = NULL;// dinvA and dwork are workspace in ztrsm
magmaDoubleComplex **cpuAarray = NULL;
magma_zmalloc( &dinvA, invA_msize * batchCount);
magma_zmalloc( &dwork, dwork_msize * batchCount );
magma_malloc_cpu((void**) &cpuAarray, batchCount*sizeof(magmaDoubleComplex*));
/* check allocation */
if ( dA_displ == NULL || dW0_displ == NULL || dW1_displ == NULL || dW2_displ == NULL ||
dW3_displ == NULL || dW4_displ == NULL || dinvA_array == NULL || dwork_array == NULL ||
dinvA == NULL || dwork == NULL || cpuAarray == NULL ) {
magma_free(dA_displ);
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dinvA_array);
magma_free(dwork_array);
magma_free( dinvA );
magma_free( dwork );
free(cpuAarray);
magma_int_t info = MAGMA_ERR_DEVICE_ALLOC;
magma_xerbla( __func__, -(info) );
return info;
}
magmablas_zlaset_q(MagmaFull, invA_msize, batchCount, MAGMA_Z_ZERO, MAGMA_Z_ZERO, dinvA, invA_msize, queue);
magmablas_zlaset_q(MagmaFull, dwork_msize, batchCount, MAGMA_Z_ZERO, MAGMA_Z_ZERO, dwork, dwork_msize, queue);
zset_pointer(dwork_array, dwork, 1, 0, 0, dwork_msize, batchCount, queue);
zset_pointer(dinvA_array, dinvA, TRI_NB, 0, 0, invA_msize, batchCount, queue);
magma_queue_t cstream;
magmablasGetKernelStream(&cstream);
magma_int_t streamid;
const magma_int_t nbstreams=32;
magma_queue_t stream[nbstreams];
for(k=0; k<nbstreams; k++){
magma_queue_create( &stream[k] );
}
magma_getvector( batchCount, sizeof(magmaDoubleComplex*), dA_array, 1, cpuAarray, 1);
magmablasSetKernelStream(NULL);
if (uplo == MagmaUpper) {
printf("Upper side is unavailable \n");
goto fin;
}
else {
for(j = 0; j < n; j+=nb) {
ib = min(nb, n-j);
#if 1
//===============================================
// panel factorization
//===============================================
magma_zdisplace_pointers(dA_displ, dA_array, ldda, j, j, batchCount, queue);
zset_pointer(dwork_array, dwork, 1, 0, 0, dwork_msize, batchCount, queue);
zset_pointer(dinvA_array, dinvA, TRI_NB, 0, 0, invA_msize, batchCount, queue);
#if 0
arginfo = magma_zpotrf_panel_batched(
uplo, n-j, ib,
dA_displ, ldda,
dwork_array, dwork_msize,
dinvA_array, invA_msize,
dW0_displ, dW1_displ, dW2_displ,
dW3_displ, dW4_displ,
info_array, j, batchCount, myhandle);
#else
//arginfo = magma_zpotrf_rectile_batched(
arginfo = magma_zpotrf_recpanel_batched(
uplo, n-j, ib, 32,
dA_displ, ldda,
dwork_array, dwork_msize,
dinvA_array, invA_msize,
dW0_displ, dW1_displ, dW2_displ,
dW3_displ, dW4_displ,
info_array, j, batchCount, myhandle, queue);
#endif
if(arginfo != 0 ) goto fin;
//===============================================
// end of panel
//===============================================
#endif
#if 1
//real_Double_t gpu_time;
//gpu_time = magma_sync_wtime(NULL);
if( (n-j-ib) > 0){
if( (n-j-ib) > gemm_crossover)
{
//-------------------------------------------
// USE STREAM HERK
//-------------------------------------------
// since it use different stream I need to wait the panel.
// But since the code use the NULL stream everywhere,
// so I don't need it, because the NULL stream do the sync by itself
//magma_queue_sync(NULL);
/* you must know the matrix layout inorder to do it */
for(k=0; k<batchCount; k++)
{
streamid = k%nbstreams;
magmablasSetKernelStream(stream[streamid]);
// call herk, class zherk must call cpu pointer
magma_zherk(MagmaLower, MagmaNoTrans, n-j-ib, ib,
d_alpha,
(const magmaDoubleComplex*) cpuAarray[k] + j+ib+j*ldda, ldda,
d_beta,
cpuAarray[k] + j+ib+(j+ib)*ldda, ldda);
}
// need to synchronise to be sure that panel do not start before
// finishing the update at least of the next panel
// BUT no need for it as soon as the other portion of the code
// use the NULL stream which do the sync by itself
//magma_device_sync();
magmablasSetKernelStream(NULL);
}
else
{
//-------------------------------------------
// USE BATCHED GEMM(which is a HERK in fact, since it only access the lower part)
//-------------------------------------------
magma_zdisplace_pointers(dA_displ, dA_array, ldda, j+ib, j, batchCount, queue);
magma_zdisplace_pointers(dW1_displ, dA_array, ldda, j+ib, j+ib, batchCount, queue);
magmablas_zherk_batched(uplo, MagmaNoTrans, n-j-ib, ib,
d_alpha, dA_displ, ldda,
d_beta, dW1_displ, ldda,
batchCount, queue);
}
}
//gpu_time = magma_sync_wtime(NULL) - gpu_time;
//real_Double_t flops = (n-j-ib) * (n-j-ib) * ib / 1e9 * batchCount;
//real_Double_t gpu_perf = flops / gpu_time;
//printf("Rows= %d, Colum=%d, herk time = %7.2fms, Gflops= %7.2f\n", n-j-ib, ib, gpu_time*1000, gpu_perf);
#endif
}
}
fin:
magma_queue_sync(NULL);
for(k=0; k<nbstreams; k++){
magma_queue_destroy( stream[k] );
}
magmablasSetKernelStream(cstream);
#if defined(USE_CUOPT)
cublasDestroy_v2(myhandle);
#endif
magma_free(dA_displ);
magma_free(dW0_displ);
magma_free(dW1_displ);
magma_free(dW2_displ);
magma_free(dW3_displ);
magma_free(dW4_displ);
magma_free(dinvA_array);
magma_free(dwork_array);
magma_free( dinvA );
magma_free( dwork );
free(cpuAarray);
return arginfo;
}