double XFSpectrumDataSerie::Magn(const int nF0, const int nF1)
{
if(nF0 == nF1)
return dB(Magnitude(m_pTCpxSpectrum[nF0]));
double dbSum = 0.0;
for(int i = nF0; i < nF1; i++)
{
if(i < int(m_unSpectrumLength))
dbSum += Magnitude(m_pTCpxSpectrum[i]);
}
return dB(dbSum/(nF1 - nF0));
}
double XFSpectrumDataSerie::Rrr(const int nF0, const int nF1)
{
if(nF0 == nF1)
return dB(SqrMod(m_pfcpxGrr[nF0]));
double dbSum = 0.0;
for(int i = nF0; i < nF1; i++)
{
if(i < int(m_unSpectrumLength))
dbSum += SqrMod(m_pfcpxGrr[i]);
}
return dB(dbSum/(nF1 - nF0));
}
void AFTimeHistoryAnalyzer::FillAvgParameters()
{
if(m_nChnlsCount > 1)
{
double* adbAvgParameters = m_aResults[m_nChnlsCount].m_adbParameters;
double dbTotalDuration = m_aResults[0].m_adbParameters[THA_TDU];
// AVG data is in the last place of m_aadbParameters vector
adbAvgParameters[THA_AVG] = dBAverage(THA_AVG);
adbAvgParameters[THA_SEL] = dBAverage(THA_AVG) + dB(dbTotalDuration);
adbAvgParameters[THA_TDU] = dbTotalDuration;
adbAvgParameters[THA_ASL] = dBAverage(THA_ASL);
adbAvgParameters[THA_THS] = Average (THA_THS);
adbAvgParameters[THA_ACT] = Average (THA_ACT);
adbAvgParameters[THA_MPK] = dBAverage(THA_MPK);
adbAvgParameters[THA_MIM] = dBAverage(THA_MIM);
adbAvgParameters[THA_MFS] = dBAverage(THA_MFS);
adbAvgParameters[THA_MSL] = dBAverage(THA_MSL);
adbAvgParameters[THA_DUR] = Average (THA_DUR);
for(int nCh = 0; nCh < m_nChnlsCount; nCh++)
{
if(m_aResults[nCh].GetParameter(THA_IMP) > 0)
adbAvgParameters[THA_IMP] = 1.0;
}
}
}
//------------------------------------------------------------------------------
void Reverb::getParameterDisplay (VstInt32 index, char* text)
{
switch (index)
{
case kParamT60low:
float2string(T60LowValue, text, kVstMaxParamStrLen);
break;
case kParamT60high:
float2string(T60HighValue, text, kVstMaxParamStrLen);
break;
case kParamTransition:
float2string(TransitionValue, text, kVstMaxParamStrLen);
break;
case kParamWetDry:
float2string(100.0*WetDryKnob, text, kVstMaxParamStrLen);
break;
case kParamQ:
float2string(ParametricQValue, text, kVstMaxParamStrLen);
break;
case kParamGamma:
float2string(dB(ParametricGammaValue), text, kVstMaxParamStrLen);
break;
case kParamFc:
float2string(ParametricFcValue, text, kVstMaxParamStrLen);
break;
default :
*text = '\0';
break;
};
}
double AFTimeHistoryAnalyzer::dBAverage(const int nParam)
{
double dbSum = 0.0;
for(int nCh = 0; nCh < m_nChnlsCount; nCh++)
dbSum += undB(m_aResults[nCh].GetParameter(nParam));
return dB(dbSum/double(m_nChnlsCount));
}
void Compressor::setupCompressor(float wthreshold,float wratio, int srate)
{
threshold = wthreshold;
ratio = wratio;
raisedTo = 1/(ratio-1);
levelEstimator.setTau(0.05, srate);
logThreshold = dB(threshold);
}
VectorXd feedForwardNetwork::rpropTrain(MatrixXd x, MatrixXd y, int numepochs, int batchsize, double incScale, double decScale, double incScaleMax, double decScaleMin, bool verbose) {
// initialize training parameters
std::vector<MatrixXd> dW(layer_size.size()-1);
std::vector<VectorXd> dB(layer_size.size()-1);
std::vector<ArrayXXi> signDeltaW(layer_size.size()-1);
std::vector<ArrayXi> signDeltaB(layer_size.size()-1);
for(int i = 0; i < layer_size.size()-1; i++) {
dW[i].setConstant(layer_size[i+1], layer_size[i], 0.1);
dB[i].setConstant(layer_size[i+1], 0.1);
signDeltaW[i].setZero(layer_size[i+1], layer_size[i]);
signDeltaB[i].setZero(layer_size[i+1]);
}
long n_sample = x.cols();
if (batchsize > n_sample) batchsize = n_sample;
int n_batch = n_sample / batchsize; // truncated if not divided
int remainder = n_sample - n_batch*batchsize;
int n_batch2 = n_batch; // n_batch2 is the actual batch number
if (remainder > 0) n_batch2++;
int s = 0; // update iteration, total iteration = numepoch x numbatch
VectorXd loss(numepochs*n_batch2); // mean sum of square error/loss
MatrixXd error; //raw error: per sample per output dimension
error.setConstant(numepochs, n_batch2, -1);
PermutationMatrix<Dynamic, Dynamic> perm(n_sample);
MatrixXd x_perm(x);
MatrixXd y_perm(y);
for (int i = 0; i < numepochs; i++) {
if (verbose) cout << "Epoch " << i + 1 << endl;
perm.setIdentity();
random_shuffle(perm.indices().data(), perm.indices().data() + perm.indices().size());
x_perm = x_perm * perm; // col = sample, shuffle samples
y_perm = y_perm * perm;
int this_batchsize = batchsize;
for(int j = 0; j < n_sample; j +=batchsize) {
if (j >= n_sample - remainder) this_batchsize = remainder;
error = ff(x_perm.middleCols(j, this_batchsize), y_perm.middleCols(j, this_batchsize));
rprop(error, dW, dB, signDeltaW, signDeltaB, incScale, decScale, incScaleMax, decScaleMin);
if (output == "softmax") {
loss[s] = -(y_perm.middleCols(j, this_batchsize).array() * post[layer_size.size()-1].array().log()).colwise().sum().mean();
} else {
loss[s] = error.array().square().mean();
}
s++;
}
}
return loss;
}
static void
usage(void)
{
const char *msg = "\
Usage: radartool [cmd]\n\
firpwr X set firpwr (thresh to check radar sig is gone) to X (int32)\n\
rrssi X set radar rssi (start det) to X dB (u_int32)\n\
height X set threshold for pulse height to X dB (u_int32)\n\
prssi set threshold to checkif pulse is gone to X dB (u_int32)\n\
inband X set threshold to check if pulse is inband to X (0.5 dB) (u_int32)\n\
dfstime X set dfs test time to X secs\n\
en_relpwr_check X enable/disable radar relative power check (AR5413 only)\n\
relpwr X set threshold to check the relative power of radar (AR5413 only)\n\
usefir128 X en/dis using in-band pwr measurement over 128 cycles(AR5413 only)\n\
en_block_check X en/dis to block OFDM weak sig as radar det(AR5413 only)\n\
en_max_rrssi X en/dis to use max rssi instead of last rssi (AR5413 only)\n\
en_relstep X en/dis to check pulse relative step (AR5413 only)\n\
relstep X set threshold to check relative step for pulse det(AR5413 only)\n\
maxlen X set max length of radar signal(in 0.8us step) (AR5413 only)\n\
numdetects get number of radar detects\n\
getnol get NOL channel information\n\
setnol set NOL channel information\n";
fprintf(stderr, "%s", msg);
}
void testDeviceVector()
{
const int aSize = 64;
std::vector<int> hA(aSize), hB(aSize);
bolt::cl::device_vector<int> dA(aSize), dB(aSize);
for(int i=0; i<aSize; i++) {
hA[i] = hB[i] = dB[i] = dA[i] = i;
};
int hSum = std::inner_product(hA.begin(), hA.end(), hB.begin(), 1);
int sum = bolt::cl::inner_product( dA.begin(), dA.end(),
dB.begin(), 1, bolt::cl::plus<int>(), bolt::cl::multiplies<int>() );
};
void Compressor::process(float input, float&output)
{
float levelEstimate,gain;
levelEstimator.process(input, levelEstimate);
float logLevel = dB(levelEstimate);
if (logLevel < logThreshold)
{
output = input;
}
else
{
float dbGain = (logLevel - logThreshold) / (1/ratio -1);
output = input * dB2lin(dbGain);
}
}
void Plot::Plot_data(vec ydata, int plot_number){
double x[ydata.length()];
double y[ydata.length()];
if(plot_number==1){
vec shifted_ydata(ydata.length());
shifted_ydata.zeros();
shifted_ydata=concat(ydata.get(ydata.length()/2,ydata.length()-1),ydata.get(0,ydata.length()/2-1));
for (int i=0;i<ydata.length();i++){
x[i]=i+1;
y[i]= dB(1.0e-14+shifted_ydata.get(i));
}
curve.setData(x,y,ydata.length());
gui->qwtPlot->replot();
}
}
/**
Purpose
-------
DSYGST reduces a real symmetric-definite generalized
eigenproblem to standard form.
If ITYPE = 1, the problem is A*x = lambda*B*x,
and A is overwritten by inv(U^H)*A*inv(U) or inv(L)*A*inv(L^H)
If ITYPE = 2 or 3, the problem is A*B*x = lambda*x or
B*A*x = lambda*x, and A is overwritten by U*A*U^H or L^H*A*L.
B must have been previously factorized as U^H*U or L*L^H by DPOTRF.
Arguments
---------
@param[in]
itype INTEGER
= 1: compute inv(U^H)*A*inv(U) or inv(L)*A*inv(L^H);
= 2 or 3: compute U*A*U^H or L^H*A*L.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored and B is factored as U^H*U;
- = MagmaLower: Lower triangle of A is stored and B is factored as L*L^H.
@param[in]
n INTEGER
The order of the matrices A and B. N >= 0.
@param[in,out]
A DOUBLE PRECISION array, dimension (LDA,N)
On entry, the symmetric matrix A. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
\n
On exit, if INFO = 0, the transformed matrix, stored in the
same format as A.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[in,out]
B DOUBLE PRECISION array, dimension (LDB,N)
The triangular factor from the Cholesky factorization of B,
as returned by DPOTRF.
B is modified by the routine but restored on exit (in lapack dsygst/dsygs2).
@param[in]
ldb INTEGER
The leading dimension of the array B. LDB >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_dsyev_comp
********************************************************************/
extern "C" magma_int_t
magma_dsygst(
magma_int_t itype, magma_uplo_t uplo, magma_int_t n,
double *A, magma_int_t lda,
double *B, magma_int_t ldb,
magma_int_t *info)
{
#define A(i_, j_) (A + (i_) + (j_)*lda)
#define B(i_, j_) (B + (i_) + (j_)*ldb)
#define dA(i_, j_) (dwork + (i_) + (j_)*ldda )
#define dB(i_, j_) (dwork + (i_) + (j_)*lddb + n*ldda)
/* Constants */
const double c_one = MAGMA_D_ONE;
const double c_neg_one = MAGMA_D_NEG_ONE;
const double c_half = MAGMA_D_HALF;
const double c_neg_half = MAGMA_D_NEG_HALF;
const double d_one = 1.0;
/* Local variables */
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t k, kb, kb2, nb;
magma_int_t ldda = n;
magma_int_t lddb = n;
magmaDouble_ptr dwork;
bool upper = (uplo == MagmaUpper);
/* Test the input parameters. */
*info = 0;
if (itype < 1 || itype > 3) {
*info = -1;
} else if (! upper && uplo != MagmaLower) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
} else if (ldb < max(1,n)) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ( n == 0 )
return *info;
if (MAGMA_SUCCESS != magma_dmalloc( &dwork, 2*n*n )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
nb = magma_get_dsygst_nb( n );
magma_queue_t queues[2];
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_create( cdev, &queues[0] );
magma_queue_create( cdev, &queues[1] );
magma_dsetmatrix( n, n, A(0, 0), lda, dA(0, 0), ldda, queues[1] );
magma_dsetmatrix( n, n, B(0, 0), ldb, dB(0, 0), lddb, queues[1] );
/* Use hybrid blocked code */
if (itype == 1) {
if (upper) {
/* Compute inv(U^H)*A*inv(U) */
for (k = 0; k < n; k += nb) {
kb = min( n-k, nb );
kb2 = min( n-k-nb, nb );
/* Update the upper triangle of A(k:n,k:n) */
lapackf77_dsygst( &itype, uplo_, &kb, A(k,k), &lda, B(k,k), &ldb, info );
magma_dsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, queues[0] );
if (k+kb < n) {
magma_dtrsm( MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
kb, n-k-kb,
c_one, dB(k,k), lddb,
dA(k,k+kb), ldda, queues[1] );
magma_queue_sync( queues[0] ); // finish set dA(k,k)
magma_dsymm( MagmaLeft, MagmaUpper,
kb, n-k-kb,
c_neg_half, dA(k,k), ldda,
dB(k,k+kb), lddb,
c_one, dA(k,k+kb), ldda, queues[1] );
magma_dsyr2k( MagmaUpper, MagmaConjTrans,
n-k-kb, kb,
c_neg_one, dA(k,k+kb), ldda,
dB(k,k+kb), lddb,
d_one, dA(k+kb,k+kb), ldda, queues[1] );
// Start copying next A block
magma_queue_sync( queues[1] );
magma_dgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(k+kb, k+kb), lda, queues[0] );
magma_dsymm( MagmaLeft, MagmaUpper,
kb, n-k-kb,
c_neg_half, dA(k,k), ldda,
dB(k,k+kb), lddb,
c_one, dA(k,k+kb), ldda, queues[1] );
magma_dtrsm( MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
kb, n-k-kb,
c_one, dB(k+kb,k+kb), lddb,
dA(k,k+kb), ldda, queues[1] );
magma_queue_sync( queues[0] ); // finish get A(k+kb,k+kb)
}
}
}
else {
/* Compute inv(L)*A*inv(L^H) */
for (k = 0; k < n; k += nb) {
kb = min( n-k, nb );
kb2 = min( n-k-nb, nb );
/* Update the lower triangle of A(k:n,k:n) */
lapackf77_dsygst( &itype, uplo_, &kb, A(k,k), &lda, B(k,k), &ldb, info );
magma_dsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, queues[0] );
if (k+kb < n) {
magma_dtrsm( MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
n-k-kb, kb,
c_one, dB(k,k), lddb,
dA(k+kb,k), ldda, queues[1] );
magma_queue_sync( queues[0] ); // finish set dA(k,k)
magma_dsymm( MagmaRight, MagmaLower,
n-k-kb, kb,
c_neg_half, dA(k,k), ldda,
dB(k+kb,k), lddb,
c_one, dA(k+kb, k), ldda, queues[1] );
magma_dsyr2k( MagmaLower, MagmaNoTrans,
n-k-kb, kb,
c_neg_one, dA(k+kb,k), ldda,
dB(k+kb,k), lddb,
d_one, dA(k+kb,k+kb), ldda, queues[1] );
// Start copying next A block
magma_queue_sync( queues[1] );
magma_dgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(k+kb, k+kb), lda, queues[0] );
magma_dsymm( MagmaRight, MagmaLower,
n-k-kb, kb,
c_neg_half, dA(k,k), ldda,
dB(k+kb,k), lddb,
c_one, dA(k+kb, k), ldda, queues[1] );
magma_dtrsm( MagmaLeft, MagmaLower, MagmaNoTrans, MagmaNonUnit,
n-k-kb, kb,
c_one, dB(k+kb,k+kb), lddb,
dA(k+kb,k), ldda, queues[1] );
magma_queue_sync( queues[0] ); // finish get A(k+kb,k+kb)
}
}
}
}
else { // itype == 2 or 3
if (upper) {
/* Compute U*A*U^H */
for (k = 0; k < n; k += nb) {
kb = min( n-k, nb );
magma_dgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(k, k), lda, queues[0] );
/* Update the upper triangle of A(1:k+kb-1,1:k+kb-1) */
if (k > 0) {
magma_dtrmm( MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
k, kb,
c_one, dB(0,0), lddb,
dA(0,k), ldda, queues[1] );
magma_dsymm( MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0,k), ldda, queues[1] );
magma_dsyr2k( MagmaUpper, MagmaNoTrans,
k, kb,
c_one, dA(0,k), ldda,
dB(0,k), lddb,
d_one, dA(0,0), ldda, queues[1] );
magma_dsymm( MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0,k), ldda, queues[1] );
magma_dtrmm( MagmaRight, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
k, kb,
c_one, dB(k,k), lddb,
dA(0,k), ldda, queues[1] );
}
magma_queue_sync( queues[0] ); // finish get A(k,k)
lapackf77_dsygst( &itype, uplo_, &kb, A(k, k), &lda, B(k, k), &ldb, info );
// this could be done on a 3rd queue
magma_dsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, queues[1] );
}
}
else {
/* Compute L^H*A*L */
for (k = 0; k < n; k += nb) {
kb = min( n-k, nb );
magma_dgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(k, k), lda, queues[0] );
/* Update the lower triangle of A(1:k+kb-1,1:k+kb-1) */
if (k > 0) {
magma_dtrmm( MagmaRight, MagmaLower, MagmaNoTrans, MagmaNonUnit,
kb, k,
c_one, dB(0,0), lddb,
dA(k,0), ldda, queues[1] );
magma_dsymm( MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k, 0), ldda, queues[1] );
magma_dsyr2k( MagmaLower, MagmaConjTrans,
k, kb,
c_one, dA(k,0), ldda,
dB(k,0), lddb,
d_one, dA(0,0), ldda, queues[1] );
magma_dsymm( MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k, 0), ldda, queues[1] );
magma_dtrmm( MagmaLeft, MagmaLower, MagmaConjTrans, MagmaNonUnit,
kb, k,
c_one, dB(k,k), lddb,
dA(k,0), ldda, queues[1] );
}
magma_queue_sync( queues[0] ); // finish get A(k,k)
lapackf77_dsygst( &itype, uplo_, &kb, A(k,k), &lda, B(k,k), &ldb, info );
// this could be done on a 3rd queue
magma_dsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, queues[1] );
}
}
}
magma_queue_sync( queues[0] ); // finish set dA(k,k) for itype 1
magma_dgetmatrix( n, n, dA(0, 0), ldda, A(0, 0), lda, queues[1] );
magma_queue_destroy( queues[0] );
magma_queue_destroy( queues[1] );
magma_free( dwork );
return *info;
} /* magma_dsygst_gpu */
void magmablas_zher2k_mgpu2(
char uplo, char trans, magma_int_t n, magma_int_t k,
magmaDoubleComplex alpha, magmaDoubleComplex *dA[], magma_int_t lda, magma_int_t aoffset,
magmaDoubleComplex *dB[], magma_int_t ldb, magma_int_t boffset,
double beta, magmaDoubleComplex *dC[], magma_int_t ldc, magma_int_t coffset,
magma_int_t ngpu, magma_int_t nb, magma_queue_t streams[][20], magma_int_t nstream )
{
#define dA(dev, i, j) (dA[dev] + (i) + (j)*lda + (aoffset) )
#define dB(dev, i, j) (dB[dev] + (i) + (j)*ldb + (boffset) )
#define dC(dev, i, j) (dC[dev] + (i) + (j)*ldc)
/* Check arguments */
magma_int_t info = 0;
if ( ! (uplo == 'l' || uplo == 'L')) {
info = -1; // 'u' not yet handled
} else if ( ! (trans == 'n' || trans == 'N')) {
info = -2; // 'c' not yet handled
} else if ( n < 0 ) {
info = -3;
} else if ( k < 0 ) {
info = -4;
} else if ( ((trans == 'n' || trans == 'N') && lda < max(1,n)) ||
((trans == 'c' || trans == 'C') && lda < max(1,k)) ) {
info = -7;
} else if ( aoffset < 0 || aoffset > lda ) {
info = -8;
} else if ( ((trans == 'n' || trans == 'N') && ldb < max(1,n)) ||
((trans == 'c' || trans == 'C') && ldb < max(1,k)) ) {
info = -10;
} else if ( boffset < 0 || boffset > ldb ) {
info = -11;
} else if ( ldc < max(1,n) ) {
info = -13;
} else if ( coffset < 0 || coffset > ldc ) {
info = -14;
} else if ( ngpu <= 0 ) {
info = -15;
} else if ( nb <= 0 ) {
info = -16;
} else if ( nstream <= 0 ) {
info = -18;
}
if ( info != 0 ) {
magma_xerbla( __func__, -(info) );
return;
}
const magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex cbeta = MAGMA_Z_MAKE( beta, 0. );
magma_int_t ib, ioff, iblock, idev, di, s;
magma_device_t cdev;
magma_queue_t cqueue;
magma_getdevice( &cdev );
magmablasGetKernelStream( &cqueue );
// loop over all blocks
// Faster to have two loops: first loop does C_hat = alpha*A*B' + beta*C
// blockoffset is offset within first block; for subsequent blocks it is 0
magma_int_t blockoffset = coffset % nb;
for( magma_int_t i = 0; i < n; i += ib ) {
ib = min( nb-blockoffset, n-i ); // block size
ioff = i + coffset; // global index in parent matrix
iblock = (ioff / nb) / ngpu; // local block id
idev = (ioff / nb) % ngpu; // device with this block
di = iblock*nb + blockoffset; // local index in parent matrix
magma_setdevice( idev );
s = iblock % nstream;
magmablasSetKernelStream( streams[ idev ][ s ] );
// C[i:n,i] = alpha * A[i:n,0] * B[i,0]' + beta*C[i:n,i]
//printf( "zgemm n=%4d, ib=%4d, k=%4d, i=%4d\n", n-i, ib, k, i );
magma_zgemm( MagmaNoTrans, MagmaConjTrans, n-i, ib, k,
alpha, dA(idev,i,0), lda,
dB(idev,i,0), ldb,
cbeta, dC(idev,ioff,di), ldc );
blockoffset = 0;
}
// second loop does C = conj(alpha)*B*A' + C_hat
alpha = MAGMA_Z_CNJG( alpha );
blockoffset = coffset % nb;
for( magma_int_t i = 0; i < n; i += ib ) {
ib = min( nb-blockoffset, n-i ); // block size
ioff = i + coffset; // global index in parent matrix
iblock = (ioff / nb) / ngpu; // local block id
idev = (ioff / nb) % ngpu; // device with this block
di = iblock*nb + blockoffset; // local index in parent matrix
magma_setdevice( idev );
s = iblock % nstream;
magmablasSetKernelStream( streams[ idev ][ s ] );
// C[i:n,i] += conj(alpha) * B[i:n,0] * A[i,0]'
//printf( "zgemm n=%4d, ib=%4d, k=%4d, i=%4d\n", n-i, ib, k, i );
magma_zgemm( MagmaNoTrans, MagmaConjTrans, n-i, ib, k,
alpha, dB(idev,i,0), ldb,
dA(idev,i,0), lda,
c_one, dC(idev,ioff,di), ldc );
blockoffset = 0;
}
magma_setdevice( cdev );
magmablasSetKernelStream( cqueue );
}
/**
Purpose
-------
ZCGESV computes the solution to a complex system of linear equations
A * X = B, A**T * X = B, or A**H * X = B,
where A is an N-by-N matrix and X and B are N-by-NRHS matrices.
ZCGESV first attempts to factorize the matrix in complex SINGLE PRECISION
and use this factorization within an iterative refinement procedure
to produce a solution with complex DOUBLE PRECISION norm-wise backward error
quality (see below). If the approach fails the method switches to a
complex DOUBLE PRECISION factorization and solve.
The iterative refinement is not going to be a winning strategy if
the ratio complex SINGLE PRECISION performance over complex DOUBLE PRECISION
performance is too small. A reasonable strategy should take the
number of right-hand sides and the size of the matrix into account.
This might be done with a call to ILAENV in the future. Up to now, we
always try iterative refinement.
The iterative refinement process is stopped if
ITER > ITERMAX
or for all the RHS we have:
RNRM < SQRT(N)*XNRM*ANRM*EPS*BWDMAX
where
o ITER is the number of the current iteration in the iterative
refinement process
o RNRM is the infinity-norm of the residual
o XNRM is the infinity-norm of the solution
o ANRM is the infinity-operator-norm of the matrix A
o EPS is the machine epsilon returned by DLAMCH('Epsilon')
The value ITERMAX and BWDMAX are fixed to 30 and 1.0D+00 respectively.
Arguments
---------
@param[in]
trans magma_trans_t
Specifies the form of the system of equations:
- = MagmaNoTrans: A * X = B (No transpose)
- = MagmaTrans: A**T * X = B (Transpose)
- = MagmaConjTrans: A**H * X = B (Conjugate transpose)
@param[in]
n INTEGER
The number of linear equations, i.e., the order of the
matrix A. N >= 0.
@param[in]
nrhs INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
@param[in,out]
dA COMPLEX_16 array on the GPU, dimension (ldda,N)
On entry, the N-by-N coefficient matrix A.
On exit, if iterative refinement has been successfully used
(info.EQ.0 and ITER.GE.0, see description below), A is
unchanged. If double precision factorization has been used
(info.EQ.0 and ITER.LT.0, see description below), then the
array dA contains 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 the array dA. ldda >= max(1,N).
@param[out]
ipiv INTEGER array, dimension (N)
The pivot indices that define the permutation matrix P;
row i of the matrix was interchanged with row IPIV(i).
Corresponds either to the single precision factorization
(if info.EQ.0 and ITER.GE.0) or the double precision
factorization (if info.EQ.0 and ITER.LT.0).
@param[out]
dipiv INTEGER array on the GPU, dimension (N)
The pivot indices; for 1 <= i <= N, after permuting, row i of the
matrix was moved to row dIPIV(i).
Note this is different than IPIV, where interchanges
are applied one-after-another.
@param[in]
dB COMPLEX_16 array on the GPU, dimension (lddb,NRHS)
The N-by-NRHS right hand side matrix B.
@param[in]
lddb INTEGER
The leading dimension of the array dB. lddb >= max(1,N).
@param[out]
dX COMPLEX_16 array on the GPU, dimension (lddx,NRHS)
If info = 0, the N-by-NRHS solution matrix X.
@param[in]
lddx INTEGER
The leading dimension of the array dX. lddx >= max(1,N).
@param
dworkd (workspace) COMPLEX_16 array on the GPU, dimension (N*NRHS)
This array is used to hold the residual vectors.
@param
dworks (workspace) COMPLEX array on the GPU, dimension (N*(N+NRHS))
This array is used to store the complex single precision matrix
and the right-hand sides or solutions in single precision.
@param[out]
iter INTEGER
- < 0: iterative refinement has failed, double precision
factorization has been performed
+ -1 : the routine fell back to full precision for
implementation- or machine-specific reasons
+ -2 : narrowing the precision induced an overflow,
the routine fell back to full precision
+ -3 : failure of SGETRF
+ -31: stop the iterative refinement after the 30th iteration
- > 0: iterative refinement has been successfully used.
Returns the number of iterations
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if info = -i, the i-th argument had an illegal value
- > 0: if info = i, U(i,i) computed in DOUBLE PRECISION is
exactly zero. The factorization has been completed,
but the factor U is exactly singular, so the solution
could not be computed.
@ingroup magma_zgesv_driver
********************************************************************/
extern "C" magma_int_t
magma_zcgesv_gpu(magma_trans_t trans, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex *dA, magma_int_t ldda,
magma_int_t *ipiv, magma_int_t *dipiv,
magmaDoubleComplex *dB, magma_int_t lddb,
magmaDoubleComplex *dX, magma_int_t lddx,
magmaDoubleComplex *dworkd, magmaFloatComplex *dworks,
magma_int_t *iter, magma_int_t *info)
{
#define dB(i,j) (dB + (i) + (j)*lddb)
#define dX(i,j) (dX + (i) + (j)*lddx)
#define dR(i,j) (dR + (i) + (j)*lddr)
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magma_int_t ione = 1;
magmaDoubleComplex *dR;
magmaFloatComplex *dSA, *dSX;
magmaDoubleComplex Xnrmv, Rnrmv;
double Anrm, Xnrm, Rnrm, cte, eps;
magma_int_t i, j, iiter, lddsa, lddr;
/* Check arguments */
*iter = 0;
*info = 0;
if ( n < 0 )
*info = -1;
else if ( nrhs < 0 )
*info = -2;
else if ( ldda < max(1,n))
*info = -4;
else if ( lddb < max(1,n))
*info = -8;
else if ( lddx < max(1,n))
*info = -10;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if ( n == 0 || nrhs == 0 )
return *info;
lddsa = n;
lddr = n;
dSA = dworks;
dSX = dSA + lddsa*n;
dR = dworkd;
eps = lapackf77_dlamch("Epsilon");
Anrm = magmablas_zlange(MagmaInfNorm, n, n, dA, ldda, (double*)dworkd );
cte = Anrm * eps * pow((double)n, 0.5) * BWDMAX;
/*
* Convert to single precision
*/
//magmablas_zlag2c( n, nrhs, dB, lddb, dSX, lddsx, info ); // done inside zcgetrs with pivots
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
magmablas_zlag2c( n, n, dA, ldda, dSA, lddsa, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// factor dSA in single precision
magma_cgetrf_gpu( n, n, dSA, lddsa, ipiv, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// Generate parallel pivots
{
magma_int_t *newipiv;
magma_imalloc_cpu( &newipiv, n );
if ( newipiv == NULL ) {
*iter = -3;
goto FALLBACK;
}
swp2pswp( trans, n, ipiv, newipiv );
magma_setvector( n, sizeof(magma_int_t), newipiv, 1, dipiv, 1 );
magma_free_cpu( newipiv );
}
// solve dSA*dSX = dB in single precision
// converts dB to dSX and applies pivots, solves, then converts result back to dX
magma_zcgetrs_gpu( trans, n, nrhs, dSA, lddsa, dipiv, dB, lddb, dX, lddx, dSX, info );
// residual dR = dB - dA*dX in double precision
magmablas_zlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dR, lddr );
if ( nrhs == 1 ) {
magma_zgemv( trans, n, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zgemm( trans, MagmaNoTrans, n, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
// TODO: use MAGMA_Z_ABS( dX(i,j) ) instead of zlange?
for( j=0; j < nrhs; j++ ) {
i = magma_izamax( n, dX(0,j), 1) - 1;
magma_zgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_izamax ( n, dR(0,j), 1 ) - 1;
magma_zgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto REFINEMENT;
}
}
*iter = 0;
return *info;
REFINEMENT:
for( iiter=1; iiter < ITERMAX; ) {
*info = 0;
// convert residual dR to single precision dSX
// solve dSA*dSX = R in single precision
// convert result back to double precision dR
// it's okay that dR is used for both dB input and dX output.
magma_zcgetrs_gpu( trans, n, nrhs, dSA, lddsa, dipiv, dR, lddr, dR, lddr, dSX, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// Add correction and setup residual
// dX += dR --and--
// dR = dB
// This saves going through dR a second time (if done with one more kernel).
// -- not really: first time is read, second time is write.
for( j=0; j < nrhs; j++ ) {
magmablas_zaxpycp( n, dR(0,j), dX(0,j), dB(0,j) );
}
// residual dR = dB - dA*dX in double precision
if ( nrhs == 1 ) {
magma_zgemv( trans, n, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zgemm( trans, MagmaNoTrans, n, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
/* Check whether the nrhs normwise backward errors satisfy the
* stopping criterion. If yes, set ITER=IITER > 0 and return. */
for( j=0; j < nrhs; j++ ) {
i = magma_izamax( n, dX(0,j), 1) - 1;
magma_zgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_izamax ( n, dR(0,j), 1 ) - 1;
magma_zgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto L20;
}
}
/* If we are here, the nrhs normwise backward errors satisfy
* the stopping criterion, we are good to exit. */
*iter = iiter;
return *info;
L20:
iiter++;
}
/* If we are at this place of the code, this is because we have
* performed ITER=ITERMAX iterations and never satisified the
* stopping criterion. Set up the ITER flag accordingly and follow
* up on double precision routine. */
*iter = -ITERMAX - 1;
FALLBACK:
/* Single-precision iterative refinement failed to converge to a
* satisfactory solution, so we resort to double precision. */
magma_zgetrf_gpu( n, n, dA, ldda, ipiv, info );
if (*info == 0) {
magmablas_zlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dX, lddx );
magma_zgetrs_gpu( trans, n, nrhs, dA, ldda, ipiv, dX, lddx, info );
}
return *info;
}
void Cpcswap_gpu( int n, cuComplex *A, int ia,int ja,int *descA, int incA,
cuComplex *B, int ib,int jb,int *descB, int incB )
{
/*
perform pcswap operation when
both distributed arrays A and B are in device memory
*/
/*
* allocate temporary space on host
* then use pcswap for communication
*/
const int use_MallocHost = FALSE;
cublasStatus cu_status;
size_t nbytes;
int elemSize = sizeof( cuComplex );
float *Atmp = 0;
float *Btmp = 0;
int descAtmp[DLEN_];
int descBtmp[DLEN_];
int ldA, ldB, ldAtmp, ldBtmp;
int nprow,npcol,myprow,mypcol;
int Locp, Locq, lrindx, lcindx, mm,nn;
int LocpA, LocqA, lrindxA, lcindxA;
int LocpB, LocqB, lrindxB, lcindxB;
int isizeA, isizeB, rsrc, csrc;
int iia,jja, iib, jjb;
int incAtmp, incBtmp;
int lrA1,lcA1, lrA2,lcA2;
int lrB1,lcB1, lrB2,lcB2;
Cblacs_gridinfo( descA[CTXT_], &nprow, &npcol, &myprow, &mypcol );
/*
* allocate storage for vector from A
*/
if (incA == 1) {
/*
* This is a column vector
*/
mm = n; nn = 1;
}
else {
/*
* This is a row vector
*/
mm = 1; nn = n;
};
setup_desc( mm,nn, ia,ja, descA, &isizeA, descAtmp );
nbytes = elemSize;
nbytes *= isizeA;
if (use_MallocHost) {
Atmp = (float *) MallocHost( nbytes );
}
else {
Atmp = (float *) malloc( nbytes );
};
assert( Atmp != 0 );
/*
* copy vector from A
*/
PROFSTART("swap:GetMatrix");
local_extent( mm,nn,ia,ja,descA, &LocpA, &LocqA, &lrA1,&lcA1, &lrA2,&lcA2 );
lrindxA = lrA1;
lcindxA = lcA1;
ldA = descA[LLD_];
ldAtmp = descAtmp[LLD_];
if ( (LocpA >= 1) && (LocqA >= 1)) {
/*
* copy from GPU device to host CPU
*/
cu_status = cublasGetMatrix( LocpA,LocqA, elemSize,
dA(lrindxA,lcindxA), ldA, Atmp, ldAtmp );
CHKERR(cu_status);
};
/*
* allocate storage for vector from B
*/
Cblacs_gridinfo( descB[CTXT_], &nprow, &npcol, &myprow, &mypcol );
if (incB == 1) {
/*
* This is a column vector
*/
mm = n; nn = 1;
}
else {
/*
* This is a row vector
*/
mm = 1; nn = n;
};
setup_desc( mm,nn, ib,jb,descB, &isizeB, descBtmp );
ldBtmp = descBtmp[LLD_];
ldB = descB[LLD_];
nbytes = elemSize;
nbytes *= isizeB;
if (use_MallocHost) {
Btmp = (float *) MallocHost( nbytes );
}
else {
Btmp = (float *) malloc( nbytes );
};
assert( Btmp != 0 );
/*
* copy vector from B
*/
local_extent( mm,nn,ib,jb,descB, &LocpB, &LocqB, &lrB1,&lcB1, &lrB2,&lcB2 );
lrindxB = lrB1;
lcindxB = lcB1;
ldB = descB[LLD_];
ldBtmp = descBtmp[LLD_];
if ((LocpB >= 1) && (LocqB >= 1)) {
/*
* Copy from GPU to CPU host
*/
cu_status = cublasGetMatrix(LocpB,LocqB,elemSize,
dB(lrindxB,lcindxB), ldB, Btmp, ldBtmp );
CHKERR(cu_status );
};
PROFEND("swap:GetMatrix");
iia = 1; jja = 1;
iib = 1; jjb = 1;
if (incA == 1) {
incAtmp = 1;
}
else {
incAtmp = descAtmp[M_];
};
if (incB == 1) {
incBtmp = 1;
}
else {
incBtmp = descBtmp[M_];
};
PROFSTART("swap:pcswap");
scalapack_pcswap( &n, Atmp, &iia, &jja, descAtmp, &incAtmp,
Btmp, &iib, &jjb, descBtmp, &incBtmp );
PROFEND("swap:pcswap");
/*
* copy from host CPU back to GPU
*/
PROFSTART("swap:SetMatrix");
if ((LocpA >= 1) && (LocqA >= 1)) {
/*
* Copy from CPU host to GPU device
*/
cu_status = cublasSetMatrix( LocpA, LocqA, elemSize,
Atmp, ldAtmp, dA(lrindxA,lcindxA), ldA );
CHKERR(cu_status);
};
if ((LocpB >= 1) && (LocqB >= 1)) {
/*
* Copy from CPU host to GPU device
*/
cu_status = cublasSetMatrix( LocpB, LocqB, elemSize,
Btmp, ldBtmp, dB(lrindxB,lcindxB), ldB );
CHKERR(cu_status);
};
PROFEND("swap:SetMatrix");
/*
* clean up
*/
if (Atmp != 0) {
if (use_MallocHost) {
FreeHost(Atmp);
}
else {
free(Atmp);
};
Atmp = 0;
};
if (Btmp != 0) {
if (use_MallocHost) {
FreeHost(Btmp);
}
else {
free(Btmp);
};
Btmp = 0;
};
return;
}
extern "C" void
magma_ssyr2k_mgpu(
magma_int_t num_gpus, magma_uplo_t uplo, magma_trans_t trans, magma_int_t nb, magma_int_t n, magma_int_t k,
float alpha,
float **db, magma_int_t lddb, magma_int_t offset_b,
float beta,
float **dc, magma_int_t lddc, magma_int_t offset,
magma_int_t num_streams, magma_queue_t stream[][10])
{
#define dB(id, i, j) (db[(id)]+(j)*lddb + (i)+offset_b)
#define dB1(id, i, j) (db[(id)]+(j)*lddb + (i)+offset_b)+k*lddb
#define dC(id, i, j) (dc[(id)]+(j)*lddc + (i))
magma_int_t i, id, ib, ii, kk, n1;
float c_one = MAGMA_S_ONE;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
/* diagonal update */
for( i=0; i < n; i += nb ) {
id = ((i+offset)/nb)%num_gpus;
kk = (i/(nb*num_gpus))%num_streams;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
ib = min(nb, n-i);
ii = nb*((i+offset)/(nb*num_gpus));
/* ssyr2k on diagonal block */
trace_gpu_start( id, kk, "syr2k", "syr2k" );
magma_ssyr2k(uplo, trans, ib, k,
alpha, dB1(id, i, 0 ), lddb,
dB(id, i, 0 ), lddb,
beta, dC(id, i+offset, ii), lddc);
trace_gpu_end( id, kk );
}
/* off-diagonal update */
if (uplo == MagmaUpper) {
for( i=nb; i < n; i += nb ) {
id = ((i+offset)/nb)%num_gpus;
kk = (i/(nb*num_gpus))%num_streams;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
ib = min(nb, n-i);
ii = nb*((i+offset)/(nb*num_gpus));
magma_sgemm(MagmaNoTrans, MagmaConjTrans, i, ib, k,
alpha, dB1(id, 0, 0 ), lddb,
dB(id, i, 0 ), lddb,
c_one, dC(id, 0, ii), lddc);
}
}
else {
for( i=0; i < n-nb; i += nb ) {
id = ((i+offset)/nb)%num_gpus;
kk = (i/(nb*num_gpus))%num_streams;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
ib = min(nb, n-i);
ii = nb*((i+offset)/(nb*num_gpus));
n1 = n-i-ib;
// sgemm on off-diagonal blocks
trace_gpu_start( id, kk, "gemm_up", "gemm_up" );
magma_sgemm(MagmaNoTrans, MagmaConjTrans, n1, ib, k,
alpha, dB1(id, i+ib, 0 ), lddb,
dB(id, i, 0 ), lddb,
c_one, dC(id, i+offset+ib, ii), lddc);
trace_gpu_end( id, kk );
}
}
if (uplo == MagmaUpper) {
for( i=nb; i < n; i += nb ) {
id = ((i+offset)/nb)%num_gpus;
kk = (i/(nb*num_gpus))%num_streams;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
ib = min(nb, n-i);
ii = nb*((i+offset)/(nb*num_gpus));
magma_sgemm(MagmaNoTrans, MagmaConjTrans, i, ib, k,
alpha, dB( id, 0, 0 ), lddb,
dB1(id, i, 0 ), lddb,
c_one, dC(id, 0, ii), lddc);
}
} else {
for( i=0; i < n-nb; i += nb ) {
id = ((i+offset)/nb)%num_gpus;
kk = (i/(nb*num_gpus))%num_streams;
magma_setdevice(id);
magmablasSetKernelStream(stream[id][kk]);
ib = min(nb, n-i);
ii = nb*((i+offset)/(nb*num_gpus));
n1 = n-i-ib;
/* sgemm on off-diagonal blocks */
trace_gpu_start( id, kk, "gemm_up", "gemm_up" );
magma_sgemm(MagmaNoTrans, MagmaConjTrans, n1, ib, k,
alpha, dB(id, i+ib, 0 ), lddb,
dB1(id, i, 0 ), lddb,
c_one, dC(id, i+offset+ib, ii), lddc);
trace_gpu_end( id, kk );
}
}
for( id=0; id < num_gpus; id++ ) {
magma_setdevice(id);
for( kk=0; kk < num_streams; kk++ ) {
magma_queue_sync(stream[id][kk]);
}
}
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
}
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, t1, t2;
double c_neg_one = MAGMA_D_NEG_ONE;
magma_int_t ione = 1;
const char trans[] = { 'N', 'C', 'T' };
const char uplo[] = { 'L', 'U' };
const char diag[] = { 'U', 'N' };
const char side[] = { 'L', 'R' };
double *A, *B, *C, *C2, *LU;
double *dA, *dB, *dC1, *dC2;
double alpha = MAGMA_D_MAKE( 0.5, 0.1 );
double beta = MAGMA_D_MAKE( 0.7, 0.2 );
double dalpha = 0.6;
double dbeta = 0.8;
double work[1], error, total_error;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t m, n, k, size, maxn, ld, info;
magma_int_t *piv;
magma_err_t err;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf( "Compares magma wrapper function to cublas function; all diffs should be exactly 0.\n\n" );
total_error = 0.;
for( int i = 0; i < opts.ntest; ++i ) {
m = opts.msize[i];
n = opts.nsize[i];
k = opts.ksize[i];
printf("=========================================================================\n");
printf( "M %d, N %d, K %d\n", (int) m, (int) n, (int) k );
// allocate matrices
// over-allocate so they can be any combination of {m,n,k} x {m,n,k}.
maxn = max( max( m, n ), k );
ld = maxn;
size = maxn*maxn;
err = magma_malloc_cpu( (void**) &piv, maxn*sizeof(magma_int_t) ); assert( err == 0 );
err = magma_dmalloc_pinned( &A, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &B, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &C, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &C2, size ); assert( err == 0 );
err = magma_dmalloc_pinned( &LU, size ); assert( err == 0 );
err = magma_dmalloc( &dA, size ); assert( err == 0 );
err = magma_dmalloc( &dB, size ); assert( err == 0 );
err = magma_dmalloc( &dC1, size ); assert( err == 0 );
err = magma_dmalloc( &dC2, size ); assert( err == 0 );
// initialize matrices
size = maxn*maxn;
lapackf77_dlarnv( &ione, ISEED, &size, A );
lapackf77_dlarnv( &ione, ISEED, &size, B );
lapackf77_dlarnv( &ione, ISEED, &size, C );
printf( "========== Level 1 BLAS ==========\n" );
// ----- test DSWAP
// swap 2nd and 3rd columns of dA, then copy to C2 and compare with A
assert( n >= 4 );
magma_dsetmatrix( m, n, A, ld, dA, ld );
magma_dsetmatrix( m, n, A, ld, dB, ld );
magma_dswap( m, dA(0,1), 1, dA(0,2), 1 );
magma_dswap( m, dB(0,1), 1, dB(0,2), 1 );
// check results, storing diff between magma and cuda calls in C2
cublasDaxpy( ld*n, c_neg_one, dA, 1, dB, 1 );
magma_dgetmatrix( m, n, dB, ld, C2, ld );
error = lapackf77_dlange( "F", &m, &k, C2, &ld, work );
total_error += error;
printf( "dswap diff %.2g\n", error );
// ----- test IDAMAX
// get argmax of column of A
magma_dsetmatrix( m, k, A, ld, dA, ld );
error = 0;
for( int j = 0; j < k; ++j ) {
magma_int_t i1 = magma_idamax( m, dA(0,j), 1 );
magma_int_t i2 = cublasIdamax( m, dA(0,j), 1 );
assert( i1 == i2 );
error += abs( i1 - i2 );
}
total_error += error;
gflops = (double)m * k / 1e9;
printf( "idamax diff %.2g\n", error );
printf( "\n" );
printf( "========== Level 2 BLAS ==========\n" );
// ----- test DGEMV
// c = alpha*A*b + beta*c, with A m*n; b,c m or n-vectors
// try no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
magma_dsetmatrix( m, n, A, ld, dA, ld );
magma_dsetvector( maxn, B, 1, dB, 1 );
magma_dsetvector( maxn, C, 1, dC1, 1 );
magma_dsetvector( maxn, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_dgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
size = (trans[ia] == 'N' ? m : n);
cublasDaxpy( size, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetvector( size, dC2, 1, C2, 1 );
error = lapackf77_dlange( "F", &size, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_DGEMV( m, n ) / 1e9;
printf( "dgemv( %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
trans[ia], error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test DSYMV
// c = alpha*A*b + beta*c, with A m*m symmetric; b,c m-vectors
// try upper/lower
for( int iu = 0; iu < 2; ++iu ) {
magma_dsetmatrix( m, m, A, ld, dA, ld );
magma_dsetvector( m, B, 1, dB, 1 );
magma_dsetvector( m, C, 1, dC1, 1 );
magma_dsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_dsymv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDsymv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( m, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_dlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYMV( m ) / 1e9;
printf( "dsymv( %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test DTRSV
// solve A*c = c, with A m*m triangular; c m-vector
// try upper/lower, no-trans/trans, unit/non-unit diag
// Factor A into LU to get well-conditioned triangles, else solve yields garbage.
// Still can give garbage if solves aren't consistent with LU factors,
// e.g., using unit diag for U, so copy lower triangle to upper triangle.
// Also used for trsm later.
lapackf77_dlacpy( "Full", &maxn, &maxn, A, &ld, LU, &ld );
lapackf77_dgetrf( &maxn, &maxn, LU, &ld, piv, &info );
for( int j = 0; j < maxn; ++j ) {
for( int i = 0; i < j; ++i ) {
*LU(i,j) = *LU(j,i);
}
}
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
magma_dsetmatrix( m, m, LU, ld, dA, ld );
magma_dsetvector( m, C, 1, dC1, 1 );
magma_dsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_dtrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDtrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( m, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_dlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_DTRSM( MagmaLeft, m, 1 ) / 1e9;
printf( "dtrsv( %c, %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], diag[id], error, gflops/t1, gflops/t2 );
}}}
printf( "\n" );
printf( "========== Level 3 BLAS ==========\n" );
// ----- test DGEMM
// C = alpha*A*B + beta*C, with A m*k or k*m; B k*n or n*k; C m*n
// try combinations of no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
for( int ib = 0; ib < 3; ++ib ) {
bool nta = (trans[ia] == 'N');
bool ntb = (trans[ib] == 'N');
magma_dsetmatrix( (nta ? m : k), (nta ? m : k), A, ld, dA, ld );
magma_dsetmatrix( (ntb ? k : n), (ntb ? n : k), B, ld, dB, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dgemm( trans[ia], trans[ib], m, n, k, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDgemm( trans[ia], trans[ib], m, n, k, alpha, dA, ld, dB, ld, beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DGEMM( m, n, k ) / 1e9;
printf( "dgemm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
trans[ia], trans[ib], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DSYMM
// C = alpha*A*B + beta*C (left) with A m*m symmetric; B,C m*n; or
// C = alpha*B*A + beta*C (right) with A n*n symmetric; B,C m*n
// try left/right, upper/lower
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
magma_dsetmatrix( m, m, A, ld, dA, ld );
magma_dsetmatrix( m, n, B, ld, dB, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dsymm( side[is], uplo[iu], m, n, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDsymm( side[is], uplo[iu], m, n, alpha, dA, ld, dB, ld, beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYMM( side[is], m, n ) / 1e9;
printf( "dsymm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
side[is], uplo[iu], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DSYRK
// C = alpha*A*A^H + beta*C (no-trans) with A m*k and C m*m symmetric; or
// C = alpha*A^H*A + beta*C (trans) with A k*m and C m*m symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
magma_dsetmatrix( n, k, A, ld, dA, ld );
magma_dsetmatrix( n, n, C, ld, dC1, ld );
magma_dsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dsyrk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDsyrk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYRK( k, n ) / 1e9;
printf( "dsyrk( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DSYR2K
// C = alpha*A*B^H + ^alpha*B*A^H + beta*C (no-trans) with A,B n*k; C n*n symmetric; or
// C = alpha*A^H*B + ^alpha*B^H*A + beta*C (trans) with A,B k*n; C n*n symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
bool nt = (trans[it] == 'N');
magma_dsetmatrix( (nt ? n : k), (nt ? n : k), A, ld, dA, ld );
magma_dsetmatrix( n, n, C, ld, dC1, ld );
magma_dsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dsyr2k( uplo[iu], trans[it], n, k, alpha, dA, ld, dB, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDsyr2k( uplo[iu], trans[it], n, k, alpha, dA, ld, dB, ld, dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DSYR2K( k, n ) / 1e9;
printf( "dsyr2k( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test DTRMM
// C = alpha*A*C (left) with A m*m triangular; C m*n; or
// C = alpha*C*A (right) with A n*n triangular; C m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == 'L');
magma_dsetmatrix( (left ? m : n), (left ? m : n), A, ld, dA, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dtrmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDtrmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DTRMM( side[is], m, n ) / 1e9;
printf( "dtrmm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// ----- test DTRSM
// solve A*X = alpha*B (left) with A m*m triangular; B m*n; or
// solve X*A = alpha*B (right) with A n*n triangular; B m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == 'L');
magma_dsetmatrix( (left ? m : n), (left ? m : n), LU, ld, dA, ld );
magma_dsetmatrix( m, n, C, ld, dC1, ld );
magma_dsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_dtrsm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasDtrsm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasDaxpy( ld*n, c_neg_one, dC1, 1, dC2, 1 );
magma_dgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_dlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_DTRSM( side[is], m, n ) / 1e9;
printf( "dtrsm( %c, %c ) diff %.2g, Gflop/s %6.2f, %6.2f\n",
uplo[iu], trans[it], error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// cleanup
magma_free_cpu( piv );
magma_free_pinned( A );
magma_free_pinned( B );
magma_free_pinned( C );
magma_free_pinned( C2 );
magma_free_pinned( LU );
magma_free( dA );
magma_free( dB );
magma_free( dC1 );
magma_free( dC2 );
}
if ( total_error != 0. ) {
printf( "total error %.2g -- ought to be 0 -- some test failed (see above).\n",
total_error );
}
else {
printf( "all tests passed\n" );
}
TESTING_FINALIZE();
return 0;
}
extern "C" magma_int_t
magma_zcposv_gpu(char uplo, magma_int_t n, magma_int_t nrhs,
magmaDoubleComplex *dA, magma_int_t ldda,
magmaDoubleComplex *dB, magma_int_t lddb,
magmaDoubleComplex *dX, magma_int_t lddx,
magmaDoubleComplex *dworkd, magmaFloatComplex *dworks,
magma_int_t *iter, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
ZCPOSV computes the solution to a complex system of linear equations
A * X = B,
where A is an N-by-N Hermitian positive definite matrix and X and B
are N-by-NRHS matrices.
ZCPOSV first attempts to factorize the matrix in complex SINGLE PRECISION
and use this factorization within an iterative refinement procedure
to produce a solution with complex DOUBLE PRECISION norm-wise backward error
quality (see below). If the approach fails the method switches to a
complex DOUBLE PRECISION factorization and solve.
The iterative refinement is not going to be a winning strategy if
the ratio complex SINGLE PRECISION performance over complex DOUBLE PRECISION
performance is too small. A reasonable strategy should take the
number of right-hand sides and the size of the matrix into account.
This might be done with a call to ILAENV in the future. Up to now, we
always try iterative refinement.
The iterative refinement process is stopped if
ITER > ITERMAX
or for all the RHS we have:
RNRM < SQRT(N)*XNRM*ANRM*EPS*BWDMAX
where
o ITER is the number of the current iteration in the iterative
refinement process
o RNRM is the infinity-norm of the residual
o XNRM is the infinity-norm of the solution
o ANRM is the infinity-operator-norm of the matrix A
o EPS is the machine epsilon returned by DLAMCH('Epsilon')
The value ITERMAX and BWDMAX are fixed to 30 and 1.0D+00 respectively.
Arguments
=========
UPLO (input) CHARACTER
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The number of linear equations, i.e., the order of the
matrix A. N >= 0.
NRHS (input) INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
dA (input or input/output) COMPLEX_16 array on the GPU, dimension (LDDA,N)
On entry, the Hermitian matrix A. If UPLO = 'U', the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if iterative refinement has been successfully used
(INFO.EQ.0 and ITER.GE.0, see description below), then A is
unchanged, if double factorization has been used
(INFO.EQ.0 and ITER.LT.0, see description below), then the
array dA contains the factor U or L from the Cholesky
factorization A = U**T*U or A = L*L**T.
LDDA (input) INTEGER
The leading dimension of the array dA. LDDA >= max(1,N).
dB (input) COMPLEX_16 array on the GPU, dimension (LDDB,NRHS)
The N-by-NRHS right hand side matrix B.
LDDB (input) INTEGER
The leading dimension of the array dB. LDDB >= max(1,N).
dX (output) COMPLEX_16 array on the GPU, dimension (LDDX,NRHS)
If INFO = 0, the N-by-NRHS solution matrix X.
LDDX (input) INTEGER
The leading dimension of the array dX. LDDX >= max(1,N).
dworkd (workspace) COMPLEX_16 array on the GPU, dimension (N*NRHS)
This array is used to hold the residual vectors.
dworks (workspace) COMPLEX array on the GPU, dimension (N*(N+NRHS))
This array is used to store the complex single precision matrix
and the right-hand sides or solutions in single precision.
ITER (output) INTEGER
< 0: iterative refinement has failed, double precision
factorization has been performed
-1 : the routine fell back to full precision for
implementation- or machine-specific reasons
-2 : narrowing the precision induced an overflow,
the routine fell back to full precision
-3 : failure of SPOTRF
-31: stop the iterative refinement after the 30th iteration
> 0: iterative refinement has been successfully used.
Returns the number of iterations
INFO (output) 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 of (DOUBLE
PRECISION) A is not positive definite, so the
factorization could not be completed, and the solution
has not been computed.
===================================================================== */
#define dB(i,j) (dB + (i) + (j)*lddb)
#define dX(i,j) (dX + (i) + (j)*lddx)
#define dR(i,j) (dR + (i) + (j)*lddr)
#define dSX(i,j) (dSX + (i) + (j)*lddsx)
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magma_int_t ione = 1;
magmaDoubleComplex *dR;
magmaFloatComplex *dSA, *dSX;
magmaDoubleComplex Xnrmv, Rnrmv;
double Anrm, Xnrm, Rnrm, cte, eps;
magma_int_t i, j, iiter, lddsa, lddsx, lddr;
/* Check arguments */
*iter = 0;
*info = 0;
if ( n < 0 )
*info = -1;
else if ( nrhs < 0 )
*info = -2;
else if ( ldda < max(1,n))
*info = -4;
else if ( lddb < max(1,n))
*info = -7;
else if ( lddx < max(1,n))
*info = -9;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if ( n == 0 || nrhs == 0 )
return *info;
lddsa = n;
lddsx = n;
lddr = n;
dSA = dworks;
dSX = dSA + lddsa*n;
dR = dworkd;
eps = lapackf77_dlamch("Epsilon");
Anrm = magmablas_zlanhe('I', uplo, n, dA, ldda, (double*)dworkd );
cte = Anrm * eps * pow((double)n, 0.5) * BWDMAX;
/*
* Convert to single precision
*/
magmablas_zlag2c( n, nrhs, dB, lddb, dSX, lddsx, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
magmablas_zlat2c( uplo, n, dA, ldda, dSA, lddsa, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// factor dSA in single precision
magma_cpotrf_gpu( uplo, n, dSA, lddsa, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// solve dSA*dSX = dB in single precision
magma_cpotrs_gpu( uplo, n, nrhs, dSA, lddsa, dSX, lddsx, info );
// residual dR = dB - dA*dX in double precision
magmablas_clag2z( n, nrhs, dSX, lddsx, dX, lddx, info );
magmablas_zlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dR, lddr );
if ( nrhs == 1 ) {
magma_zhemv( uplo, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zhemm( MagmaLeft, uplo, n, nrhs,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
// TODO: use MAGMA_Z_ABS( dX(i,j) ) instead of zlange?
for( j=0; j < nrhs; j++ ) {
i = magma_izamax( n, dX(0,j), 1) - 1;
magma_zgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_izamax ( n, dR(0,j), 1 ) - 1;
magma_zgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto REFINEMENT;
}
}
*iter = 0;
return *info;
REFINEMENT:
for( iiter=1; iiter < ITERMAX; ) {
*info = 0;
// convert residual dR to single precision dSX
magmablas_zlag2c( n, nrhs, dR, lddr, dSX, lddsx, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// solve dSA*dSX = R in single precision
magma_cpotrs_gpu( uplo, n, nrhs, dSA, lddsa, dSX, lddsx, info );
// Add correction and setup residual
// dX += dSX [including conversion] --and--
// dR = dB
for( j=0; j < nrhs; j++ ) {
magmablas_zcaxpycp( n, dSX(0,j), dX(0,j), dB(0,j), dR(0,j) );
}
// residual dR = dB - dA*dX in double precision
if ( nrhs == 1 ) {
magma_zhemv( uplo, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_zhemm( MagmaLeft, uplo, n, nrhs,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
/* Check whether the nrhs normwise backward errors satisfy the
* stopping criterion. If yes, set ITER=IITER>0 and return. */
for( j=0; j < nrhs; j++ ) {
i = magma_izamax( n, dX(0,j), 1) - 1;
magma_zgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_zlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_izamax ( n, dR(0,j), 1 ) - 1;
magma_zgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_zlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto L20;
}
}
/* If we are here, the nrhs normwise backward errors satisfy
* the stopping criterion, we are good to exit. */
*iter = iiter;
return *info;
L20:
iiter++;
}
/* If we are at this place of the code, this is because we have
* performed ITER=ITERMAX iterations and never satisified the
* stopping criterion. Set up the ITER flag accordingly and follow
* up on double precision routine. */
*iter = -ITERMAX - 1;
FALLBACK:
/* Single-precision iterative refinement failed to converge to a
* satisfactory solution, so we resort to double precision. */
magma_zpotrf_gpu( uplo, n, dA, ldda, info );
if (*info == 0) {
magmablas_zlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dX, lddx );
magma_zpotrs_gpu( uplo, n, nrhs, dA, ldda, dX, lddx, info );
}
return *info;
}
void magmablas_cher2k_mgpu2(
magma_uplo_t uplo, magma_trans_t trans, magma_int_t n, magma_int_t k,
magmaFloatComplex alpha,
magmaFloatComplex_ptr dA[], magma_int_t ldda, magma_int_t a_offset,
magmaFloatComplex_ptr dB[], magma_int_t lddb, magma_int_t b_offset,
float beta,
magmaFloatComplex_ptr dC[], magma_int_t lddc, magma_int_t c_offset,
magma_int_t ngpu, magma_int_t nb, magma_queue_t queues[][20], magma_int_t nqueue )
{
#define dA(dev, i, j) (dA[dev] + (i) + (j)*ldda + (a_offset) )
#define dB(dev, i, j) (dB[dev] + (i) + (j)*lddb + (b_offset) )
#define dC(dev, i, j) (dC[dev] + (i) + (j)*lddc)
/* Check arguments */
magma_int_t info = 0;
if ( uplo != MagmaLower ) {
info = -1; // upper not yet handled
} else if ( trans != MagmaNoTrans ) {
info = -2; // conj not yet handled
} else if ( n < 0 ) {
info = -3;
} else if ( k < 0 ) {
info = -4;
} else if ( ((trans == MagmaNoTrans) && ldda < max(1,n)) ||
((trans == Magma_ConjTrans) && ldda < max(1,k)) ) {
info = -7;
} else if ( a_offset < 0 || a_offset > ldda ) {
info = -8;
} else if ( ((trans == MagmaNoTrans) && lddb < max(1,n)) ||
((trans == Magma_ConjTrans) && lddb < max(1,k)) ) {
info = -10;
} else if ( b_offset < 0 || b_offset > lddb ) {
info = -11;
} else if ( lddc < max(1,n) ) {
info = -13;
} else if ( c_offset < 0 || c_offset > lddc ) {
info = -14;
} else if ( ngpu <= 0 ) {
info = -15;
} else if ( nb <= 0 ) {
info = -16;
} else if ( nqueue <= 0 ) {
info = -18;
}
if ( info != 0 ) {
magma_xerbla( __func__, -(info) );
return;
}
const magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex cbeta = MAGMA_C_MAKE( beta, 0. );
magma_int_t ib, ioff, iblock, idev, di, s;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
// loop over all blocks
// Faster to have two loops: first loop does C_hat = alpha*A*B**H + beta*C
// blockoffset is offset within first block; for subsequent blocks it is 0
magma_int_t blockoffset = c_offset % nb;
for( magma_int_t i = 0; i < n; i += ib ) {
ib = min( nb-blockoffset, n-i ); // block size
ioff = i + c_offset; // global index in parent matrix
iblock = (ioff / nb) / ngpu; // local block id
idev = (ioff / nb) % ngpu; // device with this block
di = iblock*nb + blockoffset; // local index in parent matrix
magma_setdevice( idev );
s = iblock % nqueue;
// C[i:n,i] = alpha * A[i:n,0] * B[i,0]' + beta*C[i:n,i]
//printf( "cgemm n=%4d, ib=%4d, k=%4d, i=%4d\n", n-i, ib, k, i );
magma_cgemm( MagmaNoTrans, Magma_ConjTrans, n-i, ib, k,
alpha, dA(idev,i,0), ldda,
dB(idev,i,0), lddb,
cbeta, dC(idev,ioff,di), lddc, queues[idev][s] );
blockoffset = 0;
}
// second loop does C = conj(alpha)*B*A**H + C_hat
alpha = MAGMA_C_CONJ( alpha );
blockoffset = c_offset % nb;
for( magma_int_t i = 0; i < n; i += ib ) {
ib = min( nb-blockoffset, n-i ); // block size
ioff = i + c_offset; // global index in parent matrix
iblock = (ioff / nb) / ngpu; // local block id
idev = (ioff / nb) % ngpu; // device with this block
di = iblock*nb + blockoffset; // local index in parent matrix
magma_setdevice( idev );
s = iblock % nqueue;
// C[i:n,i] += conj(alpha) * B[i:n,0] * A[i,0]'
//printf( "cgemm n=%4d, ib=%4d, k=%4d, i=%4d\n", n-i, ib, k, i );
magma_cgemm( MagmaNoTrans, Magma_ConjTrans, n-i, ib, k,
alpha, dB(idev,i,0), lddb,
dA(idev,i,0), ldda,
c_one, dC(idev,ioff,di), lddc, queues[idev][s] );
blockoffset = 0;
}
magma_setdevice( orig_dev );
}
void test_kronecker_product()
{
// DM = dense matrix; SM = sparse matrix
Matrix<double, 2, 3> DM_a;
SparseMatrix<double> SM_a(2,3);
SM_a.insert(0,0) = DM_a.coeffRef(0,0) = -0.4461540300782201;
SM_a.insert(0,1) = DM_a.coeffRef(0,1) = -0.8057364375283049;
SM_a.insert(0,2) = DM_a.coeffRef(0,2) = 0.3896572459516341;
SM_a.insert(1,0) = DM_a.coeffRef(1,0) = -0.9076572187376921;
SM_a.insert(1,1) = DM_a.coeffRef(1,1) = 0.6469156566545853;
SM_a.insert(1,2) = DM_a.coeffRef(1,2) = -0.3658010398782789;
MatrixXd DM_b(3,2);
SparseMatrix<double> SM_b(3,2);
SM_b.insert(0,0) = DM_b.coeffRef(0,0) = 0.9004440976767099;
SM_b.insert(0,1) = DM_b.coeffRef(0,1) = -0.2368830858139832;
SM_b.insert(1,0) = DM_b.coeffRef(1,0) = -0.9311078389941825;
SM_b.insert(1,1) = DM_b.coeffRef(1,1) = 0.5310335762980047;
SM_b.insert(2,0) = DM_b.coeffRef(2,0) = -0.1225112806872035;
SM_b.insert(2,1) = DM_b.coeffRef(2,1) = 0.5903998022741264;
SparseMatrix<double,RowMajor> SM_row_a(SM_a), SM_row_b(SM_b);
// test DM_fixedSize = kroneckerProduct(DM_block,DM)
Matrix<double, 6, 6> DM_fix_ab = kroneckerProduct(DM_a.topLeftCorner<2,3>(),DM_b);
CALL_SUBTEST(check_kronecker_product(DM_fix_ab));
CALL_SUBTEST(check_kronecker_product(kroneckerProduct(DM_a.topLeftCorner<2,3>(),DM_b)));
for(int i=0;i<DM_fix_ab.rows();++i)
for(int j=0;j<DM_fix_ab.cols();++j)
VERIFY_IS_APPROX(kroneckerProduct(DM_a,DM_b).coeff(i,j), DM_fix_ab(i,j));
// test DM_block = kroneckerProduct(DM,DM)
MatrixXd DM_block_ab(10,15);
DM_block_ab.block<6,6>(2,5) = kroneckerProduct(DM_a,DM_b);
CALL_SUBTEST(check_kronecker_product(DM_block_ab.block<6,6>(2,5)));
// test DM = kroneckerProduct(DM,DM)
MatrixXd DM_ab = kroneckerProduct(DM_a,DM_b);
CALL_SUBTEST(check_kronecker_product(DM_ab));
CALL_SUBTEST(check_kronecker_product(kroneckerProduct(DM_a,DM_b)));
// test SM = kroneckerProduct(SM,DM)
SparseMatrix<double> SM_ab = kroneckerProduct(SM_a,DM_b);
CALL_SUBTEST(check_kronecker_product(SM_ab));
SparseMatrix<double,RowMajor> SM_ab2 = kroneckerProduct(SM_a,DM_b);
CALL_SUBTEST(check_kronecker_product(SM_ab2));
CALL_SUBTEST(check_kronecker_product(kroneckerProduct(SM_a,DM_b)));
// test SM = kroneckerProduct(DM,SM)
SM_ab.setZero();
SM_ab.insert(0,0)=37.0;
SM_ab = kroneckerProduct(DM_a,SM_b);
CALL_SUBTEST(check_kronecker_product(SM_ab));
SM_ab2.setZero();
SM_ab2.insert(0,0)=37.0;
SM_ab2 = kroneckerProduct(DM_a,SM_b);
CALL_SUBTEST(check_kronecker_product(SM_ab2));
CALL_SUBTEST(check_kronecker_product(kroneckerProduct(DM_a,SM_b)));
// test SM = kroneckerProduct(SM,SM)
SM_ab.resize(2,33);
SM_ab.insert(0,0)=37.0;
SM_ab = kroneckerProduct(SM_a,SM_b);
CALL_SUBTEST(check_kronecker_product(SM_ab));
SM_ab2.resize(5,11);
SM_ab2.insert(0,0)=37.0;
SM_ab2 = kroneckerProduct(SM_a,SM_b);
CALL_SUBTEST(check_kronecker_product(SM_ab2));
CALL_SUBTEST(check_kronecker_product(kroneckerProduct(SM_a,SM_b)));
// test SM = kroneckerProduct(SM,SM) with sparse pattern
SM_a.resize(4,5);
SM_b.resize(3,2);
SM_a.resizeNonZeros(0);
SM_b.resizeNonZeros(0);
SM_a.insert(1,0) = -0.1;
SM_a.insert(0,3) = -0.2;
SM_a.insert(2,4) = 0.3;
SM_a.finalize();
SM_b.insert(0,0) = 0.4;
SM_b.insert(2,1) = -0.5;
SM_b.finalize();
SM_ab.resize(1,1);
SM_ab.insert(0,0)=37.0;
SM_ab = kroneckerProduct(SM_a,SM_b);
CALL_SUBTEST(check_sparse_kronecker_product(SM_ab));
// test dimension of result of DM = kroneckerProduct(DM,DM)
MatrixXd DM_a2(2,1);
MatrixXd DM_b2(5,4);
MatrixXd DM_ab2 = kroneckerProduct(DM_a2,DM_b2);
CALL_SUBTEST(check_dimension(DM_ab2,2*5,1*4));
DM_a2.resize(10,9);
DM_b2.resize(4,8);
DM_ab2 = kroneckerProduct(DM_a2,DM_b2);
CALL_SUBTEST(check_dimension(DM_ab2,10*4,9*8));
for(int i = 0; i < g_repeat; i++)
{
double density = Eigen::internal::random<double>(0.01,0.5);
int ra = Eigen::internal::random<int>(1,50);
int ca = Eigen::internal::random<int>(1,50);
int rb = Eigen::internal::random<int>(1,50);
int cb = Eigen::internal::random<int>(1,50);
SparseMatrix<float,ColMajor> sA(ra,ca), sB(rb,cb), sC;
SparseMatrix<float,RowMajor> sC2;
MatrixXf dA(ra,ca), dB(rb,cb), dC;
initSparse(density, dA, sA);
initSparse(density, dB, sB);
sC = kroneckerProduct(sA,sB);
dC = kroneckerProduct(dA,dB);
VERIFY_IS_APPROX(MatrixXf(sC),dC);
sC = kroneckerProduct(sA.transpose(),sB);
dC = kroneckerProduct(dA.transpose(),dB);
VERIFY_IS_APPROX(MatrixXf(sC),dC);
sC = kroneckerProduct(sA.transpose(),sB.transpose());
dC = kroneckerProduct(dA.transpose(),dB.transpose());
VERIFY_IS_APPROX(MatrixXf(sC),dC);
sC = kroneckerProduct(sA,sB.transpose());
dC = kroneckerProduct(dA,dB.transpose());
VERIFY_IS_APPROX(MatrixXf(sC),dC);
sC2 = kroneckerProduct(sA,sB);
dC = kroneckerProduct(dA,dB);
VERIFY_IS_APPROX(MatrixXf(sC2),dC);
}
}
void Cpsgecopy_general_async(int m, int n,
void *A, int ia, int ja, int *descA,
void *B, int ib, int jb, int *descB, int is_device_to_host)
{
#define dA(i,j) (((float*)A) + IDX2F(i,j,descA[LLD_]))
#define dT(i,j) (((float *)T) + IDX2F(i,j,descT[LLD_]))
#define dB(i,j) (((float *)B) + IDX2F(i,j,descB[LLD_]))
/*
perform copy
B( ib:(ib+m-1), jb:(jb+n-1)) <- A( ia:(ia+m-1),ja:(ja+n-1))
*/
const int use_MallocHost = FALSE;
const int use_igsum2d = FALSE;
cublasStatus cu_status;
cudaError_t cuda_status;
char notrans[] = "NoTrans";
int descT[DLEN_];
int ldA,ldB,ldT;
int is_same_context, is_same_mb, is_same_nb;
int is_same_p, is_same_q;
int is_same_offset;
int is_same_Locp, is_same_Locq;
int is_aligned;
int lrA1,lcA1, lrA2,lcA2;
int lrT1,lcT1, lrT2,lcT2;
int lrB1,lcB1, lrB2,lcB2;
int rsrc,csrc;
int rsrcA1,csrcA1, rsrcA2, csrcA2;
int rsrcB1,csrcB1, rsrcB2, csrcB2;
int iia,jja, iib,jjb;
int icontxt, nprow,npcol, myprow,mypcol;
int LocpA,LocqA, LocpB,LocqB, LocpT,LocqT;
int mm,nn, lmm,lnn;
size_t nbytes;
float one_[REAL_PART+IMAG_PART+1];
float *one = &(one_[0]);
float zero_[REAL_PART+IMAG_PART+1];
float *zero = &(zero_[0]);
float alpha_[REAL_PART+IMAG_PART+1];
float *alpha = &(alpha_[0]);
float beta_[REAL_PART+IMAG_PART+1];
float *beta = &(beta_[0]);
int isize, isizeT;
float *T = 0;
int elemSize = sizeof(float);
int nnb, jstart,jend,jsize;
int is_ok;
int nmax;
const int bufsize = 1024*1024;
const int use_simple = FALSE;;
one[REAL_PART] = 1.0;
one[IMAG_PART] = 0.0;
zero[REAL_PART] = 0.0;
zero[IMAG_PART] = 0.0;
if ((m <= 0) || (n <= 0)) {
return;
};
T = 0;
ldA = descA[LLD_];
ldB = descB[LLD_];
icontxt = descA[CTXT_];
Cblacs_gridinfo( icontxt, &nprow,&npcol, &myprow, &mypcol);
assert( nprow >= 1);
assert( npcol >= 1);
assert( (0 <= myprow) && (myprow < nprow));
assert( (0 <= mypcol) && (mypcol < npcol));
is_ok = (1 <= ia) && (ia + m-1 <= descA[M_]);
if (!is_ok) {
printf("Cpsgecopy (%d,%d) :ia %d m %d descA[M_] %d \n",
myprow,mypcol, ia, m, descA[M_] );
printf("Cpsgecopy (%d,%d) :ja %d n %d descA[N_] %d \n",
myprow,mypcol, ja, n, descA[N_] );
printf("Cpsgecopy (%d,%d) :ib %d jb %d descB[M_] %d descB[N_] %d\n",
myprow,mypcol, ib, jb, descB[M_], descB[N_] );
};
assert( (1 <= ia) && (ia + m-1 <= descA[M_]));
assert( (1 <= ja) && (ja + n-1 <= descA[N_]));
assert( (1 <= ib) && (ib + m-1 <= descB[M_]));
assert( (1 <= jb) && (jb + n-1 <= descB[N_]));
is_same_context = (descA[CTXT_] == descB[CTXT_]);
is_same_mb = (descA[MB_] == descB[MB_]);
is_same_nb = (descA[NB_] == descB[NB_]);
is_same_p = (Cindxg2p(ia,descA[MB_], myprow, descA[RSRC_],nprow) ==
Cindxg2p(ib,descB[MB_], myprow, descB[RSRC_],nprow) );
is_same_q = (Cindxg2p(ja,descA[NB_], mypcol, descA[CSRC_],npcol) ==
Cindxg2p(jb,descB[NB_], mypcol, descB[CSRC_],npcol) );
is_same_offset = (MOD(ia,descA[MB_]) == MOD(ib,descB[MB_])) &&
(MOD(ja,descA[NB_]) == MOD(jb,descB[NB_]));
local_extent( m,n, ia,ja,descA, &LocpA,&LocqA, &lrA1,&lcA1, &lrA2,&lcA2 );
local_extent( m,n, ib,jb,descB, &LocpB,&LocqB,&lrB1,&lcB1, &lrB2,&lcB2 );
/*
if ((LocpA >= 1) || (LocpB >= 1)) {
is_same_Locp = (LocpA == LocpB);
};
if ((LocqA >= 1) || (LocqB >= 1)) {
is_same_Locq = (LocqA == LocqB);
};
*/
is_same_Locq = (LocqA == LocqB);
is_same_Locp = (LocpA == LocpB);
is_aligned = is_same_context &&
is_same_mb && is_same_nb &&
is_same_p && is_same_q &&
is_same_offset &&
is_same_Locp && is_same_Locq;
assert( is_same_q );
assert( is_same_p );
assert( is_same_offset );
assert( is_same_Locp );
assert( is_same_Locq );
assert( is_aligned );
/*
no communication required
copy from device to host
*/
ldA = descA[LLD_];
ldB = descB[LLD_];
mm = LocpA;
nn = LocqA;
if (is_device_to_host) {
/*
* transfer from device to host
*/
if ( (mm >= 1) && (nn >= 1) ) {
#ifdef USE_CUBLASV2
{
cublasStatus_t istatus;
istatus = cublasGetMatrixAsync(mm, nn, elemSize,
(void *) dA(lrA1,lcA1), ldA, (void *) dB(lrB1,lcB1), ldB,
cublas_get_stream() );
assert( istatus == CUBLAS_STATUS_SUCCESS );
}
#else
cu_status = cublasGetMatrix(mm,nn, elemSize,
(void *) dA(lrA1,lcA1), ldA, (void *) dB(lrB1,lcB1),ldB );
CHKERR(cu_status);
#endif
};
}
else {
/*
* transfer from host to device
*/
if ( (mm >= 1) && (nn >= 1) ) {
#ifdef USE_CUBLASV2
{
cublasStatus_t istatus;
istatus = cublasSetMatrixAsync(mm,nn,elemSize,
(void *) dA(lrA1,lcA1), ldA, (void *) dB(lrB1,lcB1),ldB,
cublas_get_stream() );
assert( istatus == CUBLAS_STATUS_SUCCESS );
}
#else
cu_status = cublasSetMatrix(mm,nn,elemSize,
(void *) dA(lrA1,lcA1), ldA, (void *) dB(lrB1,lcB1),ldB );
CHKERR(cu_status);
#endif
};
};
return;
}
/**
Purpose
-------
CHEGST_GPU reduces a complex Hermitian-definite generalized
eigenproblem to standard form.
If ITYPE = 1, the problem is A*x = lambda*B*x,
and A is overwritten by inv(U**H)*A*inv(U) or inv(L)*A*inv(L**H)
If ITYPE = 2 or 3, the problem is A*B*x = lambda*x or
B*A*x = lambda*x, and A is overwritten by U*A*U**H or L**H*A*L.
B must have been previously factorized as U**H*U or L*L**H by CPOTRF.
Arguments
---------
@param[in]
itype INTEGER
= 1: compute inv(U**H)*A*inv(U) or inv(L)*A*inv(L**H);
= 2 or 3: compute U*A*U**H or L**H*A*L.
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored and B is factored as
U**H*U;
- = MagmaLower: Lower triangle of A is stored and B is factored as
L*L**H.
@param[in]
n INTEGER
The order of the matrices A and B. N >= 0.
@param[in,out]
dA COMPLEX array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = MagmaUpper, the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = MagmaLower, the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
\n
On exit, if INFO = 0, the transformed matrix, stored in the
same format as A.
@param[in]
ldda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[in]
dB COMPLEX array, dimension (LDB,N)
The triangular factor from the Cholesky factorization of B,
as returned by CPOTRF.
@param[in]
lddb INTEGER
The leading dimension of the array B. LDB >= max(1,N).
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
@ingroup magma_cheev_comp
********************************************************************/
extern "C" magma_int_t
magma_chegst_gpu(magma_int_t itype, magma_uplo_t uplo, magma_int_t n,
magmaFloatComplex *dA, magma_int_t ldda,
magmaFloatComplex *dB, magma_int_t lddb, magma_int_t *info)
{
#define A(i, j) (w + (j)*lda + (i))
#define B(i, j) (w + nb*lda + (j)*ldb + (i))
#define dA(i, j) (dA + (j)*ldda + (i))
#define dB(i, j) (dB + (j)*lddb + (i))
const char* uplo_ = lapack_uplo_const( uplo );
magma_int_t nb;
magma_int_t k, kb, kb2;
magmaFloatComplex c_one = MAGMA_C_ONE;
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex c_half = MAGMA_C_HALF;
magmaFloatComplex c_neg_half = MAGMA_C_NEG_HALF;
magmaFloatComplex *w;
magma_int_t lda;
magma_int_t ldb;
float d_one = 1.0;
int upper = (uplo == MagmaUpper);
/* Test the input parameters. */
*info = 0;
if (itype < 1 || itype > 3) {
*info = -1;
} else if (! upper && uplo != MagmaLower) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (ldda < max(1,n)) {
*info = -5;
} else if (lddb < max(1,n)) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ( n == 0 )
return *info;
nb = magma_get_chegst_nb(n);
lda = nb;
ldb = nb;
if (MAGMA_SUCCESS != magma_cmalloc_pinned( &w, 2*nb*nb )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_queue_t stream[3];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
magma_queue_create( &stream[2] );
/* Use hybrid blocked code */
if (itype == 1) {
if (upper) {
kb = min(n,nb);
/* Compute inv(U')*A*inv(U) */
magma_cgetmatrix_async( kb, kb,
dB(0, 0), lddb,
B(0, 0), nb, stream[2] );
magma_cgetmatrix_async( kb, kb,
dA(0, 0), ldda,
A(0, 0), nb, stream[1] );
for (k = 0; k < n; k += nb) {
kb = min(n-k,nb);
kb2= min(n-k-nb,nb);
/* Update the upper triangle of A(k:n,k:n) */
magma_queue_sync( stream[2] );
magma_queue_sync( stream[1] );
lapackf77_chegst( &itype, uplo_, &kb, A(0,0), &lda, B(0,0), &ldb, info);
magma_csetmatrix_async( kb, kb,
A(0, 0), lda,
dA(k, k), ldda, stream[0] );
if (k+kb < n) {
// Start copying the new B block
magma_cgetmatrix_async( kb2, kb2,
dB(k+kb, k+kb), lddb,
B(0, 0), nb, stream[2] );
magma_ctrsm(MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
kb, n-k-kb,
c_one, dB(k,k), lddb,
dA(k,k+kb), ldda);
magma_queue_sync( stream[0] );
magma_chemm(MagmaLeft, MagmaUpper,
kb, n-k-kb,
c_neg_half, dA(k,k), ldda,
dB(k,k+kb), lddb,
c_one, dA(k, k+kb), ldda);
magma_cher2k(MagmaUpper, MagmaConjTrans,
n-k-kb, kb,
c_neg_one, dA(k,k+kb), ldda,
dB(k,k+kb), lddb,
d_one, dA(k+kb,k+kb), ldda);
magma_cgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(0, 0), lda, stream[1] );
magma_chemm(MagmaLeft, MagmaUpper,
kb, n-k-kb,
c_neg_half, dA(k,k), ldda,
dB(k,k+kb), lddb,
c_one, dA(k, k+kb), ldda);
magma_ctrsm(MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
kb, n-k-kb,
c_one, dB(k+kb,k+kb), lddb,
dA(k,k+kb), ldda);
}
}
magma_queue_sync( stream[0] );
}
else {
kb = min(n,nb);
/* Compute inv(L)*A*inv(L') */
magma_cgetmatrix_async( kb, kb,
dB(0, 0), lddb,
B(0, 0), nb, stream[2] );
magma_cgetmatrix_async( kb, kb,
dA(0, 0), ldda,
A(0, 0), nb, stream[1] );
for (k = 0; k < n; k += nb) {
kb= min(n-k,nb);
kb2= min(n-k-nb,nb);
/* Update the lower triangle of A(k:n,k:n) */
magma_queue_sync( stream[2] );
magma_queue_sync( stream[1] );
lapackf77_chegst( &itype, uplo_, &kb, A(0, 0), &lda, B(0, 0), &ldb, info);
magma_csetmatrix_async( kb, kb,
A(0, 0), lda,
dA(k, k), ldda, stream[0] );
if (k+kb < n) {
// Start copying the new B block
magma_cgetmatrix_async( kb2, kb2,
dB(k+kb, k+kb), lddb,
B(0, 0), nb, stream[2] );
magma_ctrsm(MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
n-k-kb, kb,
c_one, dB(k,k), lddb,
dA(k+kb,k), ldda);
magma_queue_sync( stream[0] );
magma_chemm(MagmaRight, MagmaLower,
n-k-kb, kb,
c_neg_half, dA(k,k), ldda,
dB(k+kb,k), lddb,
c_one, dA(k+kb, k), ldda);
magma_cher2k(MagmaLower, MagmaNoTrans,
n-k-kb, kb,
c_neg_one, dA(k+kb,k), ldda,
dB(k+kb,k), lddb,
d_one, dA(k+kb,k+kb), ldda);
magma_cgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(0, 0), lda, stream[1] );
magma_chemm(MagmaRight, MagmaLower,
n-k-kb, kb,
c_neg_half, dA(k,k), ldda,
dB(k+kb,k), lddb,
c_one, dA(k+kb, k), ldda);
magma_ctrsm(MagmaLeft, MagmaLower, MagmaNoTrans, MagmaNonUnit,
n-k-kb, kb,
c_one, dB(k+kb,k+kb), lddb,
dA(k+kb,k), ldda);
}
}
}
magma_queue_sync( stream[0] );
}
else {
if (upper) {
/* Compute U*A*U' */
for (k = 0; k < n; k += nb) {
kb= min(n-k,nb);
magma_cgetmatrix_async( kb, kb,
dB(k, k), lddb,
B(0, 0), nb, stream[2] );
/* Update the upper triangle of A(1:k+kb-1,1:k+kb-1) */
if (k > 0) {
magma_ctrmm(MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
k, kb,
c_one, dB(0,0), lddb,
dA(0,k), ldda);
magma_chemm(MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0, k), ldda);
magma_queue_sync( stream[1] );
}
magma_cgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(0, 0), lda, stream[0] );
if (k > 0) {
magma_cher2k(MagmaUpper, MagmaNoTrans,
k, kb,
c_one, dA(0,k), ldda,
dB(0,k), lddb,
d_one, dA(0,0), ldda);
magma_chemm(MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0, k), ldda);
magma_ctrmm(MagmaRight, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
k, kb,
c_one, dB(k,k), lddb,
dA(0,k), ldda);
}
magma_queue_sync( stream[2] );
magma_queue_sync( stream[0] );
lapackf77_chegst( &itype, uplo_, &kb, A(0, 0), &lda, B(0, 0), &ldb, info);
magma_csetmatrix_async( kb, kb,
A(0, 0), lda,
dA(k, k), ldda, stream[1] );
}
magma_queue_sync( stream[1] );
}
else {
/* Compute L'*A*L */
for (k = 0; k < n; k += nb) {
kb= min(n-k,nb);
magma_cgetmatrix_async( kb, kb,
dB(k, k), lddb,
B(0, 0), nb, stream[2] );
/* Update the lower triangle of A(1:k+kb-1,1:k+kb-1) */
if (k > 0) {
magma_ctrmm(MagmaRight, MagmaLower, MagmaNoTrans, MagmaNonUnit,
kb, k,
c_one, dB(0,0), lddb,
dA(k,0), ldda);
magma_chemm(MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k, 0), ldda);
magma_queue_sync( stream[1] );
}
magma_cgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(0, 0), lda, stream[0] );
if (k > 0) {
magma_cher2k(MagmaLower, MagmaConjTrans,
k, kb,
c_one, dA(k,0), ldda,
dB(k,0), lddb,
d_one, dA(0,0), ldda);
magma_chemm(MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k, 0), ldda);
magma_ctrmm(MagmaLeft, MagmaLower, MagmaConjTrans, MagmaNonUnit,
kb, k,
c_one, dB(k,k), lddb,
dA(k,0), ldda);
}
magma_queue_sync( stream[2] );
magma_queue_sync( stream[0] );
lapackf77_chegst( &itype, uplo_, &kb, A(0, 0), &lda, B(0, 0), &ldb, info);
magma_csetmatrix_async( kb, kb,
A(0, 0), lda,
dA(k, k), ldda, stream[1] );
}
magma_queue_sync( stream[1] );
}
}
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_queue_destroy( stream[2] );
magma_free_pinned( w );
return *info;
} /* magma_chegst_gpu */
int main( int argc, char** argv )
{
TESTING_INIT();
real_Double_t gflops, t1, t2;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magma_int_t ione = 1;
magma_trans_t trans[] = { MagmaNoTrans, MagmaConjTrans, MagmaTrans };
magma_uplo_t uplo [] = { MagmaLower, MagmaUpper };
magma_diag_t diag [] = { MagmaUnit, MagmaNonUnit };
magma_side_t side [] = { MagmaLeft, MagmaRight };
magmaDoubleComplex *A, *B, *C, *C2, *LU;
magmaDoubleComplex *dA, *dB, *dC1, *dC2;
magmaDoubleComplex alpha = MAGMA_Z_MAKE( 0.5, 0.1 );
magmaDoubleComplex beta = MAGMA_Z_MAKE( 0.7, 0.2 );
double dalpha = 0.6;
double dbeta = 0.8;
double work[1], error, total_error;
magma_int_t ISEED[4] = {0,0,0,1};
magma_int_t m, n, k, size, maxn, ld, info;
magma_int_t *piv;
magma_int_t err;
magma_opts opts;
parse_opts( argc, argv, &opts );
printf( "Compares magma wrapper function to cublas function; all diffs should be exactly 0.\n\n" );
total_error = 0.;
for( int itest = 0; itest < opts.ntest; ++itest ) {
m = opts.msize[itest];
n = opts.nsize[itest];
k = opts.ksize[itest];
printf("=========================================================================\n");
printf( "m=%d, n=%d, k=%d\n", (int) m, (int) n, (int) k );
// allocate matrices
// over-allocate so they can be any combination of {m,n,k} x {m,n,k}.
maxn = max( max( m, n ), k );
ld = max( 1, maxn );
size = ld*maxn;
err = magma_malloc_cpu( (void**) &piv, maxn*sizeof(magma_int_t) ); assert( err == 0 );
err = magma_zmalloc_pinned( &A, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &B, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &C, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &C2, size ); assert( err == 0 );
err = magma_zmalloc_pinned( &LU, size ); assert( err == 0 );
err = magma_zmalloc( &dA, size ); assert( err == 0 );
err = magma_zmalloc( &dB, size ); assert( err == 0 );
err = magma_zmalloc( &dC1, size ); assert( err == 0 );
err = magma_zmalloc( &dC2, size ); assert( err == 0 );
// initialize matrices
size = maxn*maxn;
lapackf77_zlarnv( &ione, ISEED, &size, A );
lapackf77_zlarnv( &ione, ISEED, &size, B );
lapackf77_zlarnv( &ione, ISEED, &size, C );
printf( "========== Level 1 BLAS ==========\n" );
// ----- test ZSWAP
// swap columns 2 and 3 of dA, then copy to C2 and compare with A
if ( n >= 3 ) {
magma_zsetmatrix( m, n, A, ld, dA, ld );
magma_zsetmatrix( m, n, A, ld, dB, ld );
magma_zswap( m, dA(0,1), 1, dA(0,2), 1 );
magma_zswap( m, dB(0,1), 1, dB(0,2), 1 );
// check results, storing diff between magma and cuda calls in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dA, 1, dB, 1 );
magma_zgetmatrix( m, n, dB, ld, C2, ld );
error = lapackf77_zlange( "F", &m, &k, C2, &ld, work );
total_error += error;
printf( "zswap diff %.2g\n", error );
}
else {
printf( "zswap skipped for n < 3\n" );
}
// ----- test IZAMAX
// get argmax of column of A
magma_zsetmatrix( m, k, A, ld, dA, ld );
error = 0;
for( int j = 0; j < k; ++j ) {
magma_int_t i1 = magma_izamax( m, dA(0,j), 1 );
int i2; // NOT magma_int_t, for cublas
cublasIzamax( handle, m, dA(0,j), 1, &i2 );
// todo need sync here?
assert( i1 == i2 );
error += abs( i1 - i2 );
}
total_error += error;
gflops = (double)m * k / 1e9;
printf( "izamax diff %.2g\n", error );
printf( "\n" );
printf( "========== Level 2 BLAS ==========\n" );
// ----- test ZGEMV
// c = alpha*A*b + beta*c, with A m*n; b,c m or n-vectors
// try no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
magma_zsetmatrix( m, n, A, ld, dA, ld );
magma_zsetvector( maxn, B, 1, dB, 1 );
magma_zsetvector( maxn, C, 1, dC1, 1 );
magma_zsetvector( maxn, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_zgemv( trans[ia], m, n, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZgemv( handle, cublas_trans_const(trans[ia]),
m, n, &alpha, dA, ld, dB, 1, &beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
size = (trans[ia] == MagmaNoTrans ? m : n);
cublasZaxpy( handle, size, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetvector( size, dC2, 1, C2, 1 );
error = lapackf77_zlange( "F", &size, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZGEMV( m, n ) / 1e9;
printf( "zgemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_trans_const(trans[ia]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test ZHEMV
// c = alpha*A*b + beta*c, with A m*m symmetric; b,c m-vectors
// try upper/lower
for( int iu = 0; iu < 2; ++iu ) {
magma_zsetmatrix( m, m, A, ld, dA, ld );
magma_zsetvector( m, B, 1, dB, 1 );
magma_zsetvector( m, C, 1, dC1, 1 );
magma_zsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_zhemv( uplo[iu], m, alpha, dA, ld, dB, 1, beta, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZhemv( handle, cublas_uplo_const(uplo[iu]),
m, &alpha, dA, ld, dB, 1, &beta, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_zlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHEMV( m ) / 1e9;
printf( "zhemv( %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), error, gflops/t1, gflops/t2 );
}
printf( "\n" );
// ----- test ZTRSV
// solve A*c = c, with A m*m triangular; c m-vector
// try upper/lower, no-trans/trans, unit/non-unit diag
// Factor A into LU to get well-conditioned triangles, else solve yields garbage.
// Still can give garbage if solves aren't consistent with LU factors,
// e.g., using unit diag for U, so copy lower triangle to upper triangle.
// Also used for trsm later.
lapackf77_zlacpy( "Full", &maxn, &maxn, A, &ld, LU, &ld );
lapackf77_zgetrf( &maxn, &maxn, LU, &ld, piv, &info );
for( int j = 0; j < maxn; ++j ) {
for( int i = 0; i < j; ++i ) {
*LU(i,j) = *LU(j,i);
}
}
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
magma_zsetmatrix( m, m, LU, ld, dA, ld );
magma_zsetvector( m, C, 1, dC1, 1 );
magma_zsetvector( m, C, 1, dC2, 1 );
t1 = magma_sync_wtime( 0 );
magma_ztrsv( uplo[iu], trans[it], diag[id], m, dA, ld, dC1, 1 );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZtrsv( handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
cublas_diag_const(diag[id]), m, dA, ld, dC2, 1 );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, m, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetvector( m, dC2, 1, C2, 1 );
error = lapackf77_zlange( "F", &m, &ione, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZTRSM( MagmaLeft, m, 1 ) / 1e9;
printf( "ztrsv( %c, %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]), lapacke_diag_const(diag[id]),
error, gflops/t1, gflops/t2 );
}}}
printf( "\n" );
printf( "========== Level 3 BLAS ==========\n" );
// ----- test ZGEMM
// C = alpha*A*B + beta*C, with A m*k or k*m; B k*n or n*k; C m*n
// try combinations of no-trans/trans
for( int ia = 0; ia < 3; ++ia ) {
for( int ib = 0; ib < 3; ++ib ) {
bool nta = (trans[ia] == MagmaNoTrans);
bool ntb = (trans[ib] == MagmaNoTrans);
magma_zsetmatrix( (nta ? m : k), (nta ? m : k), A, ld, dA, ld );
magma_zsetmatrix( (ntb ? k : n), (ntb ? n : k), B, ld, dB, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zgemm( trans[ia], trans[ib], m, n, k, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZgemm( handle, cublas_trans_const(trans[ia]), cublas_trans_const(trans[ib]),
m, n, k, &alpha, dA, ld, dB, ld, &beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZGEMM( m, n, k ) / 1e9;
printf( "zgemm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_trans_const(trans[ia]), lapacke_trans_const(trans[ib]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZHEMM
// C = alpha*A*B + beta*C (left) with A m*m symmetric; B,C m*n; or
// C = alpha*B*A + beta*C (right) with A n*n symmetric; B,C m*n
// try left/right, upper/lower
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
magma_zsetmatrix( m, m, A, ld, dA, ld );
magma_zsetmatrix( m, n, B, ld, dB, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zhemm( side[is], uplo[iu], m, n, alpha, dA, ld, dB, ld, beta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZhemm( handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
m, n, &alpha, dA, ld, dB, ld, &beta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &m, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHEMM( side[is], m, n ) / 1e9;
printf( "zhemm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_side_const(side[is]), lapacke_uplo_const(uplo[iu]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZHERK
// C = alpha*A*A^H + beta*C (no-trans) with A m*k and C m*m symmetric; or
// C = alpha*A^H*A + beta*C (trans) with A k*m and C m*m symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
magma_zsetmatrix( n, k, A, ld, dA, ld );
magma_zsetmatrix( n, n, C, ld, dC1, ld );
magma_zsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zherk( uplo[iu], trans[it], n, k, dalpha, dA, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZherk( handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
n, k, &dalpha, dA, ld, &dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHERK( k, n ) / 1e9;
printf( "zherk( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZHER2K
// C = alpha*A*B^H + ^alpha*B*A^H + beta*C (no-trans) with A,B n*k; C n*n symmetric; or
// C = alpha*A^H*B + ^alpha*B^H*A + beta*C (trans) with A,B k*n; C n*n symmetric
// try upper/lower, no-trans/trans
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
bool nt = (trans[it] == MagmaNoTrans);
magma_zsetmatrix( (nt ? n : k), (nt ? n : k), A, ld, dA, ld );
magma_zsetmatrix( n, n, C, ld, dC1, ld );
magma_zsetmatrix( n, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_zher2k( uplo[iu], trans[it], n, k, alpha, dA, ld, dB, ld, dbeta, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZher2k( handle, cublas_uplo_const(uplo[iu]), cublas_trans_const(trans[it]),
n, k, &alpha, dA, ld, dB, ld, &dbeta, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( n, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZHER2K( k, n ) / 1e9;
printf( "zher2k( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}
printf( "\n" );
// ----- test ZTRMM
// C = alpha*A*C (left) with A m*m triangular; C m*n; or
// C = alpha*C*A (right) with A n*n triangular; C m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == MagmaLeft);
magma_zsetmatrix( (left ? m : n), (left ? m : n), A, ld, dA, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ztrmm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
// note cublas does trmm out-of-place (i.e., adds output matrix C),
// but allows C=B to do in-place.
t2 = magma_sync_wtime( 0 );
cublasZtrmm( handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
cublas_trans_const(trans[it]), cublas_diag_const(diag[id]),
m, n, &alpha, dA, ld, dC2, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZTRMM( side[is], m, n ) / 1e9;
printf( "ztrmm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// ----- test ZTRSM
// solve A*X = alpha*B (left) with A m*m triangular; B m*n; or
// solve X*A = alpha*B (right) with A n*n triangular; B m*n
// try left/right, upper/lower, no-trans/trans, unit/non-unit
for( int is = 0; is < 2; ++is ) {
for( int iu = 0; iu < 2; ++iu ) {
for( int it = 0; it < 3; ++it ) {
for( int id = 0; id < 2; ++id ) {
bool left = (side[is] == MagmaLeft);
magma_zsetmatrix( (left ? m : n), (left ? m : n), LU, ld, dA, ld );
magma_zsetmatrix( m, n, C, ld, dC1, ld );
magma_zsetmatrix( m, n, C, ld, dC2, ld );
t1 = magma_sync_wtime( 0 );
magma_ztrsm( side[is], uplo[iu], trans[it], diag[id], m, n, alpha, dA, ld, dC1, ld );
t1 = magma_sync_wtime( 0 ) - t1;
t2 = magma_sync_wtime( 0 );
cublasZtrsm( handle, cublas_side_const(side[is]), cublas_uplo_const(uplo[iu]),
cublas_trans_const(trans[it]), cublas_diag_const(diag[id]),
m, n, &alpha, dA, ld, dC2, ld );
t2 = magma_sync_wtime( 0 ) - t2;
// check results, storing diff between magma and cuda call in C2
cublasZaxpy( handle, ld*n, &c_neg_one, dC1, 1, dC2, 1 );
magma_zgetmatrix( m, n, dC2, ld, C2, ld );
error = lapackf77_zlange( "F", &n, &n, C2, &ld, work );
total_error += error;
gflops = FLOPS_ZTRSM( side[is], m, n ) / 1e9;
printf( "ztrsm( %c, %c ) diff %.2g, Gflop/s %7.2f, %7.2f\n",
lapacke_uplo_const(uplo[iu]), lapacke_trans_const(trans[it]),
error, gflops/t1, gflops/t2 );
}}}}
printf( "\n" );
// cleanup
magma_free_cpu( piv );
magma_free_pinned( A );
magma_free_pinned( B );
magma_free_pinned( C );
magma_free_pinned( C2 );
magma_free_pinned( LU );
magma_free( dA );
magma_free( dB );
magma_free( dC1 );
magma_free( dC2 );
fflush( stdout );
}
if ( total_error != 0. ) {
printf( "total error %.2g -- ought to be 0 -- some test failed (see above).\n",
total_error );
}
else {
printf( "all tests passed\n" );
}
TESTING_FINALIZE();
int status = (total_error != 0.);
return status;
}
/* correction_dBs corrects the dB after converted, but before scaling.*/
short dB_s(int noisefloor, float amplitude, float correction_dBs)
{
float db = dB(amplitude) + correction_dBs;
return CLAMP((int)(128.f*(db+noisefloor))/noisefloor, 0, 127);
}
extern "C" magma_int_t
magma_zhegst(magma_int_t itype, char uplo, magma_int_t n,
magmaDoubleComplex *a, magma_int_t lda,
magmaDoubleComplex *b, magma_int_t ldb, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
ZHEGST reduces a complex Hermitian-definite generalized
eigenproblem to standard form.
If ITYPE = 1, the problem is A*x = lambda*B*x,
and A is overwritten by inv(U**H)*A*inv(U) or inv(L)*A*inv(L**H)
If ITYPE = 2 or 3, the problem is A*B*x = lambda*x or
B*A*x = lambda*x, and A is overwritten by U*A*U**H or L**H*A*L.
B must have been previously factorized as U**H*U or L*L**H by ZPOTRF.
Arguments
=========
ITYPE (input) INTEGER
= 1: compute inv(U**H)*A*inv(U) or inv(L)*A*inv(L**H);
= 2 or 3: compute U*A*U**H or L**H*A*L.
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored and B is factored as
U**H*U;
= 'L': Lower triangle of A is stored and B is factored as
L*L**H.
N (input) INTEGER
The order of the matrices A and B. N >= 0.
A (input/output) COMPLEX_16 array, dimension (LDA,N)
On entry, the Hermitian matrix A. If UPLO = 'U', the leading
N-by-N upper triangular part of A contains the upper
triangular part of the matrix A, and the strictly lower
triangular part of A is not referenced. If UPLO = 'L', the
leading N-by-N lower triangular part of A contains the lower
triangular part of the matrix A, and the strictly upper
triangular part of A is not referenced.
On exit, if INFO = 0, the transformed matrix, stored in the
same format as A.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
B (input) COMPLEX_16 array, dimension (LDB,N)
The triangular factor from the Cholesky factorization of B,
as returned by ZPOTRF.
LDB (input) INTEGER
The leading dimension of the array B. LDB >= max(1,N).
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
=====================================================================*/
char uplo_[2] = {uplo, 0};
magma_int_t nb;
magma_int_t k, kb, kb2;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex c_half = MAGMA_Z_HALF;
magmaDoubleComplex c_neg_half = MAGMA_Z_NEG_HALF;
magmaDoubleComplex *dw;
magma_int_t ldda = n;
magma_int_t lddb = n;
double d_one = 1.0;
int upper = lapackf77_lsame(uplo_, "U");
/* Test the input parameters. */
*info = 0;
if (itype<1 || itype>3){
*info = -1;
}else if ((! upper) && (! lapackf77_lsame(uplo_, "L"))) {
*info = -2;
} else if (n < 0) {
*info = -3;
} else if (lda < max(1,n)) {
*info = -5;
}else if (ldb < max(1,n)) {
*info = -7;
}
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
/* Quick return */
if ( n == 0 )
return *info;
if (MAGMA_SUCCESS != magma_zmalloc( &dw, 2*n*n )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
nb = magma_get_zhegst_nb(n);
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
magma_zsetmatrix( n, n, A(0, 0), lda, dA(0, 0), ldda );
magma_zsetmatrix( n, n, B(0, 0), ldb, dB(0, 0), lddb );
/* Use hybrid blocked code */
if (itype==1) {
if (upper) {
/* Compute inv(U')*A*inv(U) */
for(k = 0; k<n; k+=nb){
kb = min(n-k,nb);
kb2= min(n-k-nb,nb);
/* Update the upper triangle of A(k:n,k:n) */
lapackf77_zhegst( &itype, uplo_, &kb, A(k,k), &lda, B(k,k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[0] );
if(k+kb<n){
magma_ztrsm(MagmaLeft, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
kb, n-k-kb,
c_one, dB(k,k), lddb,
dA(k,k+kb), ldda);
magma_queue_sync( stream[0] );
magma_zhemm(MagmaLeft, MagmaUpper,
kb, n-k-kb,
c_neg_half, dA(k,k), ldda,
dB(k,k+kb), lddb,
c_one, dA(k, k+kb), ldda);
magma_zher2k(MagmaUpper, MagmaConjTrans,
n-k-kb, kb,
c_neg_one, dA(k,k+kb), ldda,
dB(k,k+kb), lddb,
d_one, dA(k+kb,k+kb), ldda);
magma_zgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(k+kb, k+kb), lda, stream[1] );
magma_zhemm(MagmaLeft, MagmaUpper,
kb, n-k-kb,
c_neg_half, dA(k,k), ldda,
dB(k,k+kb), lddb,
c_one, dA(k, k+kb), ldda);
magma_ztrsm(MagmaRight, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
kb, n-k-kb,
c_one ,dB(k+kb,k+kb), lddb,
dA(k,k+kb), ldda);
magma_queue_sync( stream[1] );
}
}
magma_queue_sync( stream[0] );
}
else {
/* Compute inv(L)*A*inv(L') */
for(k = 0; k<n; k+=nb){
kb= min(n-k,nb);
kb2= min(n-k-nb,nb);
/* Update the lower triangle of A(k:n,k:n) */
lapackf77_zhegst( &itype, uplo_, &kb, A(k,k), &lda, B(k,k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[0] );
if(k+kb<n){
magma_ztrsm(MagmaRight, MagmaLower, MagmaConjTrans, MagmaNonUnit,
n-k-kb, kb,
c_one, dB(k,k), lddb,
dA(k+kb,k), ldda);
magma_queue_sync( stream[0] );
magma_zhemm(MagmaRight, MagmaLower,
n-k-kb, kb,
c_neg_half, dA(k,k), ldda,
dB(k+kb,k), lddb,
c_one, dA(k+kb, k), ldda);
magma_zher2k(MagmaLower, MagmaNoTrans,
n-k-kb, kb,
c_neg_one, dA(k+kb,k), ldda,
dB(k+kb,k), lddb,
d_one, dA(k+kb,k+kb), ldda);
magma_zgetmatrix_async( kb2, kb2,
dA(k+kb, k+kb), ldda,
A(k+kb, k+kb), lda, stream[1] );
magma_zhemm(MagmaRight, MagmaLower,
n-k-kb, kb,
c_neg_half, dA(k,k), ldda,
dB(k+kb,k), lddb,
c_one, dA(k+kb, k), ldda);
magma_ztrsm(MagmaLeft, MagmaLower, MagmaNoTrans, MagmaNonUnit,
n-k-kb, kb,
c_one, dB(k+kb,k+kb), lddb,
dA(k+kb,k), ldda);
}
magma_queue_sync( stream[1] );
}
}
magma_queue_sync( stream[0] );
}
else {
if (upper) {
/* Compute U*A*U' */
for(k = 0; k<n; k+=nb){
kb= min(n-k,nb);
magma_zgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(k, k), lda, stream[0] );
/* Update the upper triangle of A(1:k+kb-1,1:k+kb-1) */
if(k>0){
magma_ztrmm(MagmaLeft, MagmaUpper, MagmaNoTrans, MagmaNonUnit,
k, kb,
c_one ,dB(0,0), lddb,
dA(0,k), ldda);
magma_zhemm(MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0, k), ldda);
magma_queue_sync( stream[1] );
magma_zher2k(MagmaUpper, MagmaNoTrans,
k, kb,
c_one, dA(0,k), ldda,
dB(0,k), lddb,
d_one, dA(0,0), ldda);
magma_zhemm(MagmaRight, MagmaUpper,
k, kb,
c_half, dA(k,k), ldda,
dB(0,k), lddb,
c_one, dA(0, k), ldda);
magma_ztrmm(MagmaRight, MagmaUpper, MagmaConjTrans, MagmaNonUnit,
k, kb,
c_one, dB(k,k), lddb,
dA(0,k), ldda);
}
magma_queue_sync( stream[0] );
lapackf77_zhegst( &itype, uplo_, &kb, A(k, k), &lda, B(k, k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[1] );
}
magma_queue_sync( stream[1] );
}
else {
/* Compute L'*A*L */
for(k = 0; k<n; k+=nb){
kb= min(n-k,nb);
magma_zgetmatrix_async( kb, kb,
dA(k, k), ldda,
A(k, k), lda, stream[0] );
/* Update the lower triangle of A(1:k+kb-1,1:k+kb-1) */
if(k>0){
magma_ztrmm(MagmaRight, MagmaLower, MagmaNoTrans, MagmaNonUnit,
kb, k,
c_one ,dB(0,0), lddb,
dA(k,0), ldda);
magma_zhemm(MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k, 0), ldda);
magma_queue_sync( stream[1] );
magma_zher2k(MagmaLower, MagmaConjTrans,
k, kb,
c_one, dA(k,0), ldda,
dB(k,0), lddb,
d_one, dA(0,0), ldda);
magma_zhemm(MagmaLeft, MagmaLower,
kb, k,
c_half, dA(k,k), ldda,
dB(k,0), lddb,
c_one, dA(k, 0), ldda);
magma_ztrmm(MagmaLeft, MagmaLower, MagmaConjTrans, MagmaNonUnit,
kb, k,
c_one, dB(k,k), lddb,
dA(k,0), ldda);
}
magma_queue_sync( stream[0] );
lapackf77_zhegst( &itype, uplo_, &kb, A(k,k), &lda, B(k,k), &ldb, info);
magma_zsetmatrix_async( kb, kb,
A(k, k), lda,
dA(k, k), ldda, stream[1] );
}
magma_queue_sync( stream[1] );
}
}
magma_zgetmatrix( n, n, dA(0, 0), ldda, A(0, 0), lda );
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free( dw );
return *info;
} /* magma_zhegst_gpu */
// ----------------------------------------------------------------------
extern "C" void
magma_zher2k_mgpu(
magma_int_t ngpu,
magma_uplo_t uplo, magma_trans_t trans, magma_int_t nb, magma_int_t n, magma_int_t k,
magmaDoubleComplex alpha,
magmaDoubleComplex_ptr dB[], magma_int_t lddb, magma_int_t b_offset,
double beta,
magmaDoubleComplex_ptr dC[], magma_int_t lddc, magma_int_t c_offset,
magma_int_t nqueue, magma_queue_t queues[][10])
{
#define dB(id, i, j) (dB[(id)] + (j)*lddb + (i)+b_offset)
#define dB1(id, i, j) (dB[(id)] + (j)*lddb + (i)+b_offset) + k*lddb
#define dC(id, i, j) (dC[(id)] + (j)*lddc + (i))
magma_int_t i, id, ib, ii, kk, n1;
magmaDoubleComplex c_one = MAGMA_Z_ONE;
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
/* diagonal update */
for( i=0; i < n; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = (i/(nb*ngpu))%nqueue;
magma_setdevice( id );
magmablasSetKernelStream( queues[id][kk] );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
/* zher2k on diagonal block */
trace_gpu_start( id, kk, "syr2k", "syr2k" );
magma_zher2k( uplo, trans, ib, k,
alpha, dB1(id, i, 0 ), lddb,
dB(id, i, 0 ), lddb,
beta, dC(id, i+c_offset, ii), lddc );
trace_gpu_end( id, kk );
}
/* off-diagonal update */
if (uplo == MagmaUpper) {
for( i=nb; i < n; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = (i/(nb*ngpu))%nqueue;
magma_setdevice( id );
magmablasSetKernelStream( queues[id][kk] );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
magma_zgemm( MagmaNoTrans, MagmaConjTrans, i, ib, k,
alpha, dB1(id, 0, 0 ), lddb,
dB(id, i, 0 ), lddb,
c_one, dC(id, 0, ii), lddc );
}
}
else {
for( i=0; i < n-nb; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = (i/(nb*ngpu))%nqueue;
magma_setdevice( id );
magmablasSetKernelStream( queues[id][kk] );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
n1 = n-i-ib;
// zgemm on off-diagonal blocks
trace_gpu_start( id, kk, "gemm_up", "gemm_up" );
magma_zgemm( MagmaNoTrans, MagmaConjTrans, n1, ib, k,
alpha, dB1(id, i+ib, 0 ), lddb,
dB(id, i, 0 ), lddb,
c_one, dC(id, i+c_offset+ib, ii), lddc );
trace_gpu_end( id, kk );
}
}
if (uplo == MagmaUpper) {
for( i=nb; i < n; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = (i/(nb*ngpu))%nqueue;
magma_setdevice( id );
magmablasSetKernelStream( queues[id][kk] );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
magma_zgemm( MagmaNoTrans, MagmaConjTrans, i, ib, k,
alpha, dB( id, 0, 0 ), lddb,
dB1(id, i, 0 ), lddb,
c_one, dC(id, 0, ii), lddc );
}
} else {
for( i=0; i < n-nb; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = (i/(nb*ngpu))%nqueue;
magma_setdevice( id );
magmablasSetKernelStream( queues[id][kk] );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
n1 = n-i-ib;
/* zgemm on off-diagonal blocks */
trace_gpu_start( id, kk, "gemm_up", "gemm_up" );
magma_zgemm( MagmaNoTrans, MagmaConjTrans, n1, ib, k,
alpha, dB(id, i+ib, 0 ), lddb,
dB1(id, i, 0 ), lddb,
c_one, dC(id, i+c_offset+ib, ii), lddc );
trace_gpu_end( id, kk );
}
}
for( id=0; id < ngpu; id++ ) {
magma_setdevice( id );
for( kk=0; kk < nqueue; kk++ ) {
magma_queue_sync( queues[id][kk] );
}
}
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
}
extern "C" magma_int_t
magma_dsgesv_gpu(char trans, magma_int_t n, magma_int_t nrhs,
double *dA, magma_int_t ldda,
magma_int_t *ipiv, magma_int_t *dipiv,
double *dB, magma_int_t lddb,
double *dX, magma_int_t lddx,
double *dworkd, float *dworks,
magma_int_t *iter, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
DSGESV computes the solution to a real system of linear equations
A * X = B or A' * X = B
where A is an N-by-N matrix and X and B are N-by-NRHS matrices.
DSGESV first attempts to factorize the matrix in real SINGLE PRECISION
and use this factorization within an iterative refinement procedure
to produce a solution with real DOUBLE PRECISION norm-wise backward error
quality (see below). If the approach fails the method switches to a
real DOUBLE PRECISION factorization and solve.
The iterative refinement is not going to be a winning strategy if
the ratio real SINGLE PRECISION performance over real DOUBLE PRECISION
performance is too small. A reasonable strategy should take the
number of right-hand sides and the size of the matrix into account.
This might be done with a call to ILAENV in the future. Up to now, we
always try iterative refinement.
The iterative refinement process is stopped if
ITER > ITERMAX
or for all the RHS we have:
RNRM < SQRT(N)*XNRM*ANRM*EPS*BWDMAX
where
o ITER is the number of the current iteration in the iterative
refinement process
o RNRM is the infinity-norm of the residual
o XNRM is the infinity-norm of the solution
o ANRM is the infinity-operator-norm of the matrix A
o EPS is the machine epsilon returned by DLAMCH('Epsilon')
The value ITERMAX and BWDMAX are fixed to 30 and 1.0D+00 respectively.
Arguments
=========
TRANS (input) CHARACTER*1
Specifies the form of the system of equations:
= 'N': A * X = B (No transpose)
= 'T': A'* X = B (Transpose)
= 'C': A'* X = B (Conjugate transpose = Transpose)
N (input) INTEGER
The number of linear equations, i.e., the order of the
matrix A. N >= 0.
NRHS (input) INTEGER
The number of right hand sides, i.e., the number of columns
of the matrix B. NRHS >= 0.
dA (input or input/output) DOUBLE PRECISION array on the GPU, dimension (ldda,N)
On entry, the N-by-N coefficient matrix A.
On exit, if iterative refinement has been successfully used
(info.EQ.0 and ITER.GE.0, see description below), A is
unchanged. If double precision factorization has been used
(info.EQ.0 and ITER.LT.0, see description below), then the
array dA contains the factors L and U from the factorization
A = P*L*U; the unit diagonal elements of L are not stored.
ldda (input) INTEGER
The leading dimension of the array dA. ldda >= max(1,N).
IPIV (output) INTEGER array, dimension (N)
The pivot indices that define the permutation matrix P;
row i of the matrix was interchanged with row IPIV(i).
Corresponds either to the single precision factorization
(if info.EQ.0 and ITER.GE.0) or the double precision
factorization (if info.EQ.0 and ITER.LT.0).
dIPIV (output) INTEGER array on the GPU, dimension (min(M,N))
The pivot indices; for 1 <= i <= min(M,N), row i of the
matrix was moved to row IPIV(i).
dB (input) DOUBLE PRECISION array on the GPU, dimension (lddb,NRHS)
The N-by-NRHS right hand side matrix B.
lddb (input) INTEGER
The leading dimension of the array dB. lddb >= max(1,N).
dX (output) DOUBLE PRECISION array on the GPU, dimension (lddx,NRHS)
If info = 0, the N-by-NRHS solution matrix X.
lddx (input) INTEGER
The leading dimension of the array dX. lddx >= max(1,N).
dworkd (workspace) DOUBLE PRECISION array on the GPU, dimension (N*NRHS)
This array is used to hold the residual vectors.
dworks (workspace) SINGLE PRECISION array on the GPU, dimension (N*(N+NRHS))
This array is used to store the real single precision matrix
and the right-hand sides or solutions in single precision.
iter (output) INTEGER
< 0: iterative refinement has failed, double precision
factorization has been performed
-1 : the routine fell back to full precision for
implementation- or machine-specific reasons
-2 : narrowing the precision induced an overflow,
the routine fell back to full precision
-3 : failure of SGETRF
-31: stop the iterative refinement after the 30th iteration
> 0: iterative refinement has been successfully used.
Returns the number of iterations
info (output) INTEGER
= 0: successful exit
< 0: if info = -i, the i-th argument had an illegal value
> 0: if info = i, U(i,i) computed in DOUBLE PRECISION is
exactly zero. The factorization has been completed,
but the factor U is exactly singular, so the solution
could not be computed.
===================================================================== */
#define dB(i,j) (dB + (i) + (j)*lddb)
#define dX(i,j) (dX + (i) + (j)*lddx)
#define dR(i,j) (dR + (i) + (j)*lddr)
double c_neg_one = MAGMA_D_NEG_ONE;
double c_one = MAGMA_D_ONE;
magma_int_t ione = 1;
double *dR;
float *dSA, *dSX;
double Xnrmv, Rnrmv;
double Anrm, Xnrm, Rnrm, cte, eps;
magma_int_t i, j, iiter, lddsa, lddr;
/* Check arguments */
*iter = 0;
*info = 0;
if ( n < 0 )
*info = -1;
else if ( nrhs < 0 )
*info = -2;
else if ( ldda < max(1,n))
*info = -4;
else if ( lddb < max(1,n))
*info = -8;
else if ( lddx < max(1,n))
*info = -10;
if (*info != 0) {
magma_xerbla( __func__, -(*info) );
return *info;
}
if ( n == 0 || nrhs == 0 )
return *info;
lddsa = n;
lddr = n;
dSA = dworks;
dSX = dSA + lddsa*n;
dR = dworkd;
eps = lapackf77_dlamch("Epsilon");
Anrm = magmablas_dlange('I', n, n, dA, ldda, (double*)dworkd );
cte = Anrm * eps * pow((double)n, 0.5) * BWDMAX;
/*
* Convert to single precision
*/
//magmablas_dlag2s( n, nrhs, dB, lddb, dSX, lddsx, info ); // done inside dsgetrs with pivots
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
magmablas_dlag2s( n, n, dA, ldda, dSA, lddsa, info );
if (*info != 0) {
*iter = -2;
goto FALLBACK;
}
// factor dSA in single precision
magma_sgetrf_gpu( n, n, dSA, lddsa, ipiv, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// Generate parallel pivots
{
magma_int_t *newipiv;
magma_imalloc_cpu( &newipiv, n );
if ( newipiv == NULL ) {
*iter = -3;
goto FALLBACK;
}
swp2pswp( trans, n, ipiv, newipiv );
magma_setvector( n, sizeof(magma_int_t), newipiv, 1, dipiv, 1 );
magma_free_cpu( newipiv );
}
// solve dSA*dSX = dB in single precision
// converts dB to dSX and applies pivots, solves, then converts result back to dX
magma_dsgetrs_gpu( trans, n, nrhs, dSA, lddsa, dipiv, dB, lddb, dX, lddx, dSX, info );
// residual dR = dB - dA*dX in double precision
magmablas_dlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dR, lddr );
if ( nrhs == 1 ) {
magma_dgemv( trans, n, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_dgemm( trans, MagmaNoTrans, n, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
// TODO: use MAGMA_D_ABS( dX(i,j) ) instead of dlange?
for( j=0; j < nrhs; j++ ) {
i = magma_idamax( n, dX(0,j), 1) - 1;
magma_dgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_dlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_idamax ( n, dR(0,j), 1 ) - 1;
magma_dgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_dlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto REFINEMENT;
}
}
*iter = 0;
return *info;
REFINEMENT:
for( iiter=1; iiter < ITERMAX; ) {
*info = 0;
// convert residual dR to single precision dSX
// solve dSA*dSX = R in single precision
// convert result back to double precision dR
// it's okay that dR is used for both dB input and dX output.
magma_dsgetrs_gpu( trans, n, nrhs, dSA, lddsa, dipiv, dR, lddr, dR, lddr, dSX, info );
if (*info != 0) {
*iter = -3;
goto FALLBACK;
}
// Add correction and setup residual
// dX += dR --and--
// dR = dB
// This saves going through dR a second time (if done with one more kernel).
// -- not really: first time is read, second time is write.
for( j=0; j < nrhs; j++ ) {
magmablas_daxpycp( n, dR(0,j), dX(0,j), dB(0,j) );
}
// residual dR = dB - dA*dX in double precision
if ( nrhs == 1 ) {
magma_dgemv( trans, n, n,
c_neg_one, dA, ldda,
dX, 1,
c_one, dR, 1 );
}
else {
magma_dgemm( trans, MagmaNoTrans, n, nrhs, n,
c_neg_one, dA, ldda,
dX, lddx,
c_one, dR, lddr );
}
/* Check whether the nrhs normwise backward errors satisfy the
* stopping criterion. If yes, set ITER=IITER>0 and return. */
for( j=0; j < nrhs; j++ ) {
i = magma_idamax( n, dX(0,j), 1) - 1;
magma_dgetmatrix( 1, 1, dX(i,j), 1, &Xnrmv, 1 );
Xnrm = lapackf77_dlange( "F", &ione, &ione, &Xnrmv, &ione, NULL );
i = magma_idamax ( n, dR(0,j), 1 ) - 1;
magma_dgetmatrix( 1, 1, dR(i,j), 1, &Rnrmv, 1 );
Rnrm = lapackf77_dlange( "F", &ione, &ione, &Rnrmv, &ione, NULL );
if ( Rnrm > Xnrm*cte ) {
goto L20;
}
}
/* If we are here, the nrhs normwise backward errors satisfy
* the stopping criterion, we are good to exit. */
*iter = iiter;
return *info;
L20:
iiter++;
}
/* If we are at this place of the code, this is because we have
* performed ITER=ITERMAX iterations and never satisified the
* stopping criterion. Set up the ITER flag accordingly and follow
* up on double precision routine. */
*iter = -ITERMAX - 1;
FALLBACK:
/* Single-precision iterative refinement failed to converge to a
* satisfactory solution, so we resort to double precision. */
magma_dgetrf_gpu( n, n, dA, ldda, ipiv, info );
if (*info == 0) {
magmablas_dlacpy( MagmaUpperLower, n, nrhs, dB, lddb, dX, lddx );
magma_dgetrs_gpu( trans, n, nrhs, dA, ldda, ipiv, dX, lddx, info );
}
return *info;
}
void magmablas_ssymm_mgpu_spec(
magma_side_t side, magma_uplo_t uplo, magma_int_t m, magma_int_t n,
float alpha,
magmaFloat_ptr dA[], magma_int_t ldda, magma_int_t offset,
magmaFloat_ptr dB[], magma_int_t lddb,
float beta,
magmaFloat_ptr dC[], magma_int_t lddc,
magmaFloat_ptr dwork[], magma_int_t dworksiz,
float *C, magma_int_t ldc,
float *work[], magma_int_t worksiz, // TODO unused
magma_int_t ngpu, magma_int_t nb,
magma_queue_t queues[][20], magma_int_t nqueue,
magma_event_t redevents[][MagmaMaxGPUs*MagmaMaxGPUs+10], magma_int_t nbevents,
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2], magma_int_t nbcmplx )
{
#define dA(dev, i, j) (dA[dev] + (i) + (j)*ldda)
#define dB(dev, i, j) (dB[dev] + (i) + (j)*lddb)
#define dC(dev, i, j) (dC[dev] + (i) + (j)*lddc)
#define dwork(dev, i, j) (dwork[dev] + (i) + (j)*lddwork)
#define C(i, j) (C + (i) + (j)*ldc)
if ( side != MagmaLeft || uplo != MagmaLower ) {
fprintf( stderr, "%s: only Left Lower implemented\n", __func__ );
}
assert( ldda >= m );
assert( lddb >= m );
assert( lddc >= m );
assert( nqueue >= ngpu );
assert( nbevents >= ngpu*ngpu );
magmaFloat_ptr dwork1[MagmaMaxGPUs];
magmaFloat_ptr dwork2[MagmaMaxGPUs];
magma_int_t lddwork = lddc;
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
dwork1[dev] = dwork[dev];
dwork2[dev] = dwork[dev]+n*lddwork;
}
assert( dworksiz >= (2*n*lddwork) );
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_t cstream;
magmablasGetKernelStream(&cstream);
magma_int_t dev, devperm, myblk, mycolsize, myblkoffst;
magma_int_t gdev, gcolsize, gmaster, gngpu;
magma_int_t masterdev, lcdev, lccolsize, myngpu;
magma_int_t stdev = (offset/nb)%ngpu;
magma_int_t blockoffset = offset % nb;
magma_int_t fstblksiz = 0;
if(blockoffset>0){
fstblksiz = min(m, (nb - blockoffset));
}
//magma_int_t nbblk = magma_ceildiv(m, nb);
magma_int_t nbblk = magma_ceildiv((m+blockoffset), nb);
magma_int_t maxgsize = n*nb*magma_ceildiv(nbblk, ngpu);
magma_int_t remm = m- fstblksiz;
magma_int_t nbblkoffst = offset/nb;
magma_int_t nblstblks = -1;
magma_int_t devlstblk = -1;
magma_int_t lstblksiz = remm%nb;
if(lstblksiz>0){
nblstblks = nbblk%ngpu;
devlstblk = (nblstblks-1+ngpu)%ngpu;
}
magma_int_t nbcmplxactive = 0;
magma_int_t cmplxisactive[MagmaMaxGPUs];
magma_int_t gpuisactive[MagmaMaxGPUs];
memset(gpuisactive, 0, MagmaMaxGPUs*sizeof(magma_int_t));
memset(cmplxisactive, 0, MagmaMaxGPUs*sizeof(magma_int_t));
//*******************************
// each GPU make a GEMM with the
// transpose of its blocks to compute
// a final portion of X=A*VT
//*******************************
/* dB = V*T already ==> dB**H = T**H * V**H
* compute T**H * V**H * X is equal to compute locally (VT)**H_i*X_i
* then each GPU broadcast its X_i to assemble the full X which is used
* to compute W = X - 0.5 * V * T**H * V**H * X = X - 0.5 * V *dwork3
*/
if(ngpu ==1){
magma_setdevice( 0 );
magmablasSetKernelStream( queues[ 0 ][ 0 ] );
// compute X[me] = A*VT = A[me]^tr *VT;
magma_sgemm( MagmaConjTrans, MagmaNoTrans, m, n, m,
alpha, dA(0, offset, offset), ldda,
dB[0], lddb,
beta, dC[0], lddc );
return;
}
//ngpu>1
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
masterdev = -1;
gnode[cmplxid][MagmaMaxGPUs+1] = -1;
myngpu = gnode[cmplxid][MagmaMaxGPUs];
for( magma_int_t idev = 0; idev < myngpu; ++idev ) {
dev = gnode[cmplxid][idev];
devperm = (dev-stdev+ngpu)%ngpu;
myblk = (nbblk/ngpu) + (nbblk%ngpu > devperm ? 1:0 );
mycolsize = myblk*nb;
myblkoffst = nb*((nbblkoffst/ngpu)+(nbblkoffst%ngpu > dev?1:0));
if(dev==stdev){
mycolsize -= blockoffset;
myblkoffst += blockoffset; // local index in parent matrix
}
if((devperm==devlstblk)&&(lstblksiz>0)){
mycolsize -= (nb-(remm%nb));
}
mycolsize = min(mycolsize, m);
if(mycolsize>0){
if(masterdev==-1) masterdev = dev;
//printf("dev %d devperm %d on cmplx %d master %d nbblk %d myblk %d m %d n %d mycolsize %d stdev %d fstblksize %d lastdev %d lastsize %d dA(%d, %d, %d) ==> dwork(%d, %d)\n", dev, devperm, cmplxid, masterdev, nbblk, myblk, m, n, mycolsize, stdev, fstblksiz, devlstblk, remm%nb, dev, offset, myblkoffst, dev, maxgsize*dev);
gpuisactive[dev] = mycolsize;
magma_setdevice( dev );
magmablasSetKernelStream( queues[ dev ][ dev ] );
magma_sgemm( MagmaConjTrans, MagmaNoTrans, mycolsize, n, m,
alpha, dA(dev, offset, myblkoffst), ldda,
dB(dev, 0, 0), lddb,
beta, &dwork[dev][maxgsize*dev], mycolsize );
magma_event_record(redevents[dev][dev*ngpu+dev], queues[dev][dev]);
}
if(dev == masterdev){
nbcmplxactive = nbcmplxactive +1;
cmplxisactive[cmplxid] = 1;
gnode[cmplxid][MagmaMaxGPUs+1] = masterdev;
}
}
}
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_queue_sync( queues[ dev ][ dev ] );
}
*/
//*******************************
// each Master GPU has the final
// result either by receiving
// from CPU of by making the add
// by himself, so now it is time
// to broadcast over the GPUs of
// its board.
//*******************************
//printf("=======================================================================\n");
//printf(" sending \n");
//printf("=======================================================================\n");
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
for( magma_int_t idev = 0; idev < myngpu; ++idev ) {
dev = gnode[cmplxid][idev];
mycolsize = gpuisactive[dev];
if(mycolsize>0){
// I am an active GPU send my portion local
// to all active gpu of my cmplex and global to the
// active master of the other real and they should
// send it out to their actives slaves.
magma_setdevice( dev );
//==============================================
// sending to the master of the active real
//==============================================
//printf ("\n\n**************GPU %d\n ", dev);
//printf (" GPU %d sending to cmplx masters\n", dev);
for( magma_int_t k = 0; k < nbcmplx; ++k ) {
if(k!=cmplxid){
gmaster = gnode[k][MagmaMaxGPUs+1];
if(gmaster!=-1){ //real is active
//printf (" device %d from cmplx %d is sending to master %d on cmplx %d block of size %d event %d\n", dev, cmplxid, gmaster, k, mycolsize, redevents[dev][gmaster*ngpu+dev]);
magma_queue_wait_event(queues[ dev ][ gmaster ], redevents[dev][dev*ngpu+dev]);
magma_scopymatrix_async(
mycolsize, n,
&dwork[dev ][maxgsize*dev], mycolsize,
&dwork[gmaster][maxgsize*dev], mycolsize, queues[dev][gmaster] );
magma_event_record(redevents[dev][gmaster*ngpu+dev], queues[dev][gmaster]);
}
}
}
//==============================================
//
//==============================================
// sending to the active GPUs of my real
//==============================================
//printf (" GPU %d sending internal\n", dev);
for( magma_int_t l = 0; l < myngpu; ++l ) {
lcdev = gnode[cmplxid][l];
lccolsize = gpuisactive[lcdev];
if((lcdev!=dev)&&(lccolsize>0)){
//printf (" device %d from cmplx %d is sending internal to dev %d block of size %d event %d\n", dev, cmplxid, lcdev, mycolsize, redevents[dev][lcdev*ngpu+dev]);
magma_queue_wait_event(queues[ dev ][ lcdev ], redevents[dev][dev*ngpu+dev]);
magma_scopymatrix_async(
mycolsize, n,
&dwork[dev ][maxgsize*dev], mycolsize,
&dwork[lcdev][maxgsize*dev], mycolsize, queues[dev][lcdev] );
magma_event_record(redevents[dev][lcdev*ngpu+dev], queues[dev][lcdev]);
}
}
//==============================================
}// end if mycolsize>0
}// for idev
}// for cmplxid
//printf("=======================================================================\n");
//printf(" master wait and resend internally \n");
//printf("=======================================================================\n");
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//==============================================
// if I am active master so wait receiving contribution
// of the GPUs of other real and send it locally
//==============================================
if(masterdev != -1){
mycolsize = gpuisactive[masterdev];
magma_setdevice( masterdev );
//printf(" GPU %d distributing internal\n", masterdev);
for( magma_int_t k = 0; k < nbcmplx; ++k ) {
if(k!=cmplxid){
gngpu = gnode[k][MagmaMaxGPUs];
for( magma_int_t g = 0; g < gngpu; ++g ) {
gdev = gnode[k][g];
gcolsize = gpuisactive[gdev];
// check if I received from this GPU,
// if yes send it to my group
if(gcolsize>0){
magma_queue_wait_event(queues[ masterdev ][ gdev ], redevents[gdev][masterdev*ngpu+gdev]);
for( magma_int_t l = 0; l < myngpu; ++l ) {
lcdev = gnode[cmplxid][l];
lccolsize = gpuisactive[lcdev];
if((lcdev!=masterdev)&&(lccolsize>0)){
//printf(" Master %d on cmplx %d waiting on event %d is distributing internal results of %d to lcdev %d block of size %d event %d\n", masterdev, cmplxid, redevents[gdev][masterdev*ngpu+gdev], gdev, lcdev, gcolsize, redevents[masterdev][lcdev*ngpu+gdev]);
magma_scopymatrix_async(
gcolsize, n,
&dwork[masterdev][maxgsize*gdev], gcolsize,
&dwork[lcdev ][maxgsize*gdev], gcolsize, queues[masterdev][gdev] );
magma_event_record(redevents[masterdev][lcdev*ngpu+gdev], queues[masterdev][gdev]);
}
}
}
}
}
}
}// if active master
//==============================================
}// for cmplxid
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_queue_sync( queues[ dev ][ 0 ] );
for( magma_int_t s = 0; s < ngpu; ++s ) {
magma_queue_sync( queues[ dev ][ s ] );
}
}
*/
//printf("=======================================================================\n");
//printf(" distributing \n");
//printf("=======================================================================\n");
magma_int_t lcblki, gbblki, gblk, ib;
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
for( magma_int_t idev = 0; idev < myngpu; ++idev ) {
dev = gnode[cmplxid][idev];
mycolsize = gpuisactive[dev];
if(mycolsize>0){ // I am an active GPU
//printf("\n\n==============GPU %d collecting\n", dev);
magma_setdevice( dev );
// collect my results first as tyhere is no need to wait to
// receive nothing, just wait that my gemm are done.
// in theory this should be inside the loop but cuda was not
// able to run it first for all gpu and on gpu>0 it was waiting
// however it was on different stream so it should run. but maybe
// this is because there are too many function call and this make
// cuda not handleit so nice. anyway it coul dbe removed when cuda
// is able to lunch it first without wait.
gdev = dev;
gcolsize = gpuisactive[gdev];
if(gcolsize>0){
devperm = (gdev-stdev+ngpu)%ngpu;
gblk = (nbblk/ngpu) + (nbblk%ngpu > devperm ? 1:0 );
magmablasSetKernelStream( queues[ dev ][ gdev ] );
magma_queue_wait_event(queues[ dev ][ gdev ], redevents[gdev][dev*ngpu+gdev]);
//printf (" GPU %d stream %d doing slacpy\n", dev, queues[ dev ][ gdev ]);
for( magma_int_t blki = 0; blki < gblk; ++blki){
gbblki = (blki*ngpu + devperm)*nb - blockoffset;
lcblki = blki*nb;
ib = nb;//min(nb, m-gbblki);
if(gdev==stdev){
lcblki = blki*nb-blockoffset;
if(blki==0){
gbblki = 0;
lcblki = 0;
ib = nb-blockoffset;
}
}
ib = min(ib, m-gbblki);
//printf(" blockoffset %d nbblk %d stdev %d receiving from gdev %d gblk %d gcolsize %d copying blki %d of size ibxn %dx%d from work[%d] to C[%d]\n", blockoffset, nbblk, stdev, gdev, gblk, gcolsize, blki, ib, n, lcblki, gbblki);
magmablas_slacpy( MagmaFull, ib, n, &dwork[dev][maxgsize*gdev+lcblki], gcolsize, &dC[dev][gbblki], lddc);
}// end blki
}
for( magma_int_t k = 0; k < nbcmplx; ++k ) {
gngpu = gnode[k][MagmaMaxGPUs];
for( magma_int_t g = 0; g < gngpu; ++g ) {
gdev = gnode[k][g];
gcolsize = gpuisactive[gdev];
// if gcolsize>0, ==> gpu gdev was active and so
// I received from him/computed a portion of dwork,
// so go over its gblk and distribute it on dC.
if(gdev!=dev){
if(gcolsize>0){
devperm = (gdev-stdev+ngpu)%ngpu;
gblk = (nbblk/ngpu) + (nbblk%ngpu > devperm ? 1:0 );
magmablasSetKernelStream( queues[ dev ][ gdev ] );
if(k==cmplxid){
//we are on the same group so wait on event issued by gdev for me citing his id
magma_queue_wait_event(queues[ dev ][ gdev ], redevents[gdev][dev*ngpu+gdev]);
//printf (" GPU %d queue %d waiting on event %d to collecte from %d the size of gcolsize %d\n", dev, queues[ dev ][ gdev ], redevents[gdev][dev*ngpu+gdev], gdev, gcolsize);
}else{
//we are on different group so:
//if I am the master wait on the event issued by gdev for me citing his id
//else wait event issued by my master for me on the behalf of gdev
//printf (" GPU %d queue %d waiting on event %d to collecte from %d the size of gcolsize %d\n", dev, queues[ dev ][ gdev ], redevents[masterdev][dev*ngpu+gdev], gdev, gcolsize);
if(dev==masterdev)
magma_queue_wait_event(queues[ dev ][ gdev ], redevents[gdev][dev*ngpu+gdev]);
else
magma_queue_wait_event(queues[ dev ][ gdev ], redevents[masterdev][dev*ngpu+gdev]);
}
//printf (" GPU %d stream %d doing slacpy\n", dev, queues[ dev ][ gdev ]);
for( magma_int_t blki = 0; blki < gblk; ++blki){
gbblki = (blki*ngpu + devperm)*nb - blockoffset;
lcblki = blki*nb;
ib = nb;//min(nb, m-gbblki);
if(gdev==stdev){
lcblki = blki*nb-blockoffset;
if(blki==0){
gbblki = 0;
lcblki = 0;
ib = nb-blockoffset;
}
}
ib = min(ib, m-gbblki);
//printf(" blockoffset %d nbblk %d stdev %d receiving from gdev %d gblk %d gcolsize %d copying blki %d of size ibxn %dx%d from work[%d] to C[%d]\n", blockoffset, nbblk, stdev, gdev, gblk, gcolsize, blki, ib, n, lcblki, gbblki);
magmablas_slacpy( MagmaFull, ib, n, &dwork[dev][maxgsize*gdev+lcblki], gcolsize, &dC[dev][gbblki], lddc);
}// end blki
}// en gcolsize>0 meaning gdev is active
} // end if gdev != dev
}// end loop over the g gpus of the cmplx k
}//end loop over the real k
}// end mycolsize>0 meaning that I am active
}// end loop over idev of cmplxid
}// end loop of the cmplx
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_device_sync();
}
// put back the input gpu and its input stream
magma_setdevice( cdev );
magmablasSetKernelStream( cstream );
}
void magmablas_ssymm_mgpu_com(
magma_side_t side, magma_uplo_t uplo, magma_int_t m, magma_int_t n,
float alpha,
float *dA[], magma_int_t ldda, magma_int_t offset,
float *dB[], magma_int_t lddb,
float beta, float *dC[], magma_int_t lddc,
float *dwork[], magma_int_t dworksiz,
float *C, magma_int_t ldc,
float *work[], magma_int_t worksiz,
magma_int_t ngpu, magma_int_t nb,
magma_queue_t streams[][20], magma_int_t nstream,
magma_event_t redevents[][MagmaMaxGPUs*MagmaMaxGPUs+10], magma_int_t nbevents,
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2], magma_int_t nbcmplx )
{
#define dA(dev, i, j) (dA[dev] + (i) + (j)*ldda)
#define dB(dev, i, j) (dB[dev] + (i) + (j)*lddb)
#define dC(dev, i, j) (dC[dev] + (i) + (j)*lddc)
#define dwork(dev, i, j) (dwork[dev] + (i) + (j)*lddwork)
#define C(i, j) (C + (i) + (j)*ldc)
//printf("####################################################\n");
//printf(" start ssymm \n");
//printf("####################################################\n");
if ( side != MagmaLeft || uplo != MagmaLower ) {
fprintf( stderr, "%s: only Left Lower implemented\n", __func__ );
}
assert( ldda >= m );
assert( lddb >= m );
assert( lddc >= m );
assert( nstream >= ngpu );
assert( nbevents >= ngpu*ngpu );
float c_one = MAGMA_S_ONE;
float *dwork1[MagmaMaxGPUs];
float *dwork2[MagmaMaxGPUs];
magma_int_t maxgsize = n*m;
magma_int_t lddwork = lddc;
magma_int_t ldwork = m;
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
dwork1[dev] = dwork[dev]; // size of dwork1 is n*lddwork
dwork2[dev] = dwork[dev]+n*lddwork; // size of dwork2 is maxgsize*ngpu
}
assert( dworksiz >= (n*lddwork+maxgsize*ngpu) );
assert( worksiz >= (n*ldwork) );
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_t cstream;
magmablasGetKernelStream(&cstream);
magma_int_t dev, devperm, myblk, mycolsize, myblkoffst;
magma_int_t gmaster;
magma_int_t masterdev, lcdev, lccolsize, myngpu;
magma_int_t stdev = (offset/nb)%ngpu;
magma_int_t blockoffset = offset % nb;
magma_int_t fstblksiz = 0;
if(blockoffset>0){
fstblksiz = min(m, (nb - blockoffset));
}
//magma_int_t nbblk = magma_ceildiv(m, nb);
magma_int_t nbblk = magma_ceildiv((m+blockoffset), nb);
magma_int_t remm = m- fstblksiz;
magma_int_t nbblkoffst = offset/nb;
magma_int_t nblstblks = -1;
magma_int_t devlstblk = -1;
magma_int_t lstblksiz = remm%nb;
if(lstblksiz>0){
nblstblks = nbblk%ngpu;
devlstblk = (nblstblks-1+ngpu)%ngpu;
}
magma_int_t nbcmplxactive = 0;
magma_int_t cmplxisactive[MagmaMaxGPUs];
magma_int_t gpuisactive[MagmaMaxGPUs];
memset(gpuisactive, 0, MagmaMaxGPUs*sizeof(magma_int_t));
memset(cmplxisactive, 0, MagmaMaxGPUs*sizeof(magma_int_t));
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
cudaMemset(dwork(dev,0,0), 0, (lddwork)*(n)*sizeof(float) );
// put all dC on all dev to 0 except the one which
// hold i==0 because this one has to multiply by beta.
if(dev!=stdev){
cudaMemset(dC(dev,0,0), 0, (lddc)*(n)*sizeof(float) );
}
}
magma_int_t newoffset = offset;
// 1. symmetrize
if(blockoffset>0){
newoffset = offset+fstblksiz; // newoffset is adjusted over nb
magma_int_t myblkoffst = (nbblkoffst/ngpu)+(nbblkoffst%ngpu > stdev?1:0);
//printf("STDEV %d voici offset %d remm %d myblockoffset %d siz %d \n", stdev, offset, remm, myblkoffst, fstblksiz);
magma_setdevice( stdev );
magmablasSetKernelStream( streams[ stdev ][ 0 ] );
magmablas_ssymmetrize_tiles( MagmaLower, fstblksiz, dA(stdev, offset, myblkoffst*nb+blockoffset), ldda, 1, ngpu*nb, nb );
}
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_int_t newstdev = (newoffset/nb)%ngpu;
magma_int_t nbblk = remm/nb; // number of block of size nb. if m%nb>0 then a last block exist and is of size ib=m%nb
magma_int_t myblk = (nbblk/ngpu) + (nbblk%ngpu > ((dev-newstdev+ngpu)%ngpu) ? 1:0 );
magma_int_t devperm = (dev-newstdev+ngpu)%ngpu;
magma_int_t nbblkoffst = newoffset/nb;
magma_int_t myblkoffst = (nbblkoffst/ngpu)+(nbblkoffst%ngpu > dev?1:0);
//printf("dev %d devperm %d newoffset %d rowoff %d coloff %d myblk %d \n", dev, devperm, newoffset, newoffset+devperm*nb, myblkoffst*nb, myblk);
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
magmablas_ssymmetrize_tiles( MagmaLower, nb, dA(dev, newoffset+devperm*nb, myblkoffst*nb), ldda, myblk, ngpu*nb, nb );
if(remm%nb>0){
magma_int_t nblstblks = (nbblk+1)%ngpu;
magma_int_t devlstblk = (nblstblks-1+ngpu)%ngpu;
//printf("==> siz %d devperm %d, devlstblk %d, newoffset+nbblk*nb %d, myblkoffst*nb+ myblk*nb %d\n", remm % nb, devperm, devlstblk, newoffset+nbblk*nb, myblkoffst*nb+ myblk*nb);
if(devperm==devlstblk)
magmablas_ssymmetrize( MagmaLower, remm % nb, dA(dev, newoffset+nbblk*nb, myblkoffst*nb+ myblk*nb), ldda ); // last partial tile
}
}
/*
magma_int_t siz = m+offset;
float *R;
magma_smalloc_cpu( &R, siz*siz );
// collecte back A
magmablas_sgetmatrix_1D_bcyclic( siz, siz, dA, ldda, R, siz, ngpu, nb );
magma_setdevice( 0 );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
//magma_sgetmatrix( siz, siz, dA[0], ldda, R, siz );
FILE *trace_file;
trace_file = fopen("AJETE/Aafter", "w");
for (int j = 0; j < siz ; j++)
for (int i = 0; i < siz ; i++)
fprintf(trace_file, "%10d%10d%40.30e\n", i+1, j+1, R[j*siz+i]);
fclose(trace_file);
return;
*/
// ROW GEMM transpose a row and make a gemm with a block
// if only 1 GPU used the ROW GEMM is integrated with the
// COL GEMM (better accuracy observed) and better perf
if(ngpu>1){
for( magma_int_t i = fstblksiz; i < m; i += nb ) {
magma_int_t ib = min( nb, m-i ); // block size
magma_int_t ioff = i + offset; // start global index in parent matrix
//magma_int_t dev = (ioff / nb) % ngpu;
magma_int_t nbblkoffst = offset/nb;
magma_int_t nbblk = magma_ceildiv(i, nb);
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_int_t myblk = (nbblk/ngpu) + (nbblk%ngpu > ((dev-stdev+ngpu)%ngpu) ? 1:0 );
magma_int_t myblkoffst = (nbblkoffst/ngpu)+(nbblkoffst%ngpu > dev?1:0);
magma_int_t myrowsize = myblk * nb;
magma_int_t coloffset = myblkoffst*nb;
if(dev==stdev) {
myrowsize = myrowsize -blockoffset;
coloffset = myblkoffst*nb+blockoffset;
}
//printf("ROW GEMM: voici i %d ib %d ioff %d nbblkoffst %d stdev %d dev %d myblk %d myblkoffset %d coloffset %d rowsize %d\n", i, ib, ioff, nbblkoffst, stdev, dev, myblk, myblkoffst, coloffset, myrowsize);
if(myrowsize>0){
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 1 ] );
magma_sgemm( MagmaConjTrans, MagmaNoTrans, myrowsize, n, ib,
alpha, dA(dev,ioff,coloffset), ldda,
dB(dev,i,0), lddb,
c_one, dwork(dev,0,0), lddwork );
}
}
}
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_event_record(redevents[dev][1], streams[dev][1]);
}
}
// COL GEMM
// blockoffset is offset within first block; for subsequent blocks it is 0
if(blockoffset>0){
magma_int_t ib = min( nb-blockoffset, m ); // block size
magma_int_t iblock = (offset / nb) / ngpu; // local block id
magma_int_t di = iblock*nb+blockoffset; // local index in parent matrix
magma_setdevice( stdev );
magmablasSetKernelStream( streams[ stdev ][ 0 ] );
//printf("DEV %d COL GEMM first ioff %d di %d m %d n %d ib %d \n", stdev, offset, di, m, n, ib);
magma_sgemm( MagmaNoTrans, MagmaNoTrans, m, n, ib,
alpha, dA(stdev,offset,di), ldda,
dB(stdev,0,0), lddb,
beta, dC(stdev,0,0), lddc );
}
// COL GEMM
for( magma_int_t i = fstblksiz; i < m; i += nb ) {
magma_int_t ib = min( nb, m-i ); // block size
magma_int_t ioff = i + offset; // start global index in parent matrix
magma_int_t iblock = (ioff / nb) / ngpu; // local block id
magma_int_t dev = (ioff / nb) % ngpu;
magma_int_t di = iblock*nb; // local index in parent matrix
//printf("DEV %d COL GEMM i %d ioff %d di %d m-i %d n %d ib %d \n", dev, i, ioff, di, m-i, n, ib);
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
if(i==0){
magma_sgemm( MagmaNoTrans, MagmaNoTrans, m-i, n, ib,
alpha, dA(dev,ioff,di), ldda,
dB(dev,i,0), lddb,
beta, dC(dev,i,0), lddc );
}else{
magma_sgemm( MagmaNoTrans, MagmaNoTrans, m-i, n, ib,
alpha, dA(dev,ioff,di), ldda,
dB(dev,i,0), lddb,
c_one, dC(dev,i,0), lddc );
}
magma_event_record(redevents[dev][0], streams[dev][0]);
// if only 1 GPU is used, do the ROW GEMM
if(ngpu==1){
// NOTE THAT because the COL gemm write dC below the diagonal (i)
// and the ROW GEMM write dC from 0 to diag-1, so they could
// run in parallel on different streams.
//
// NO NO NO because
// it might happen that col finished i and strated i+1 while row still at i
// magmablasSetKernelStream( streams[ dev ][ 0 ] );
magma_sgemm( MagmaConjTrans, MagmaNoTrans, i, n, ib,
alpha, dA(dev,ioff,offset), ldda,
dB(dev,i,0), lddb,
c_one, dC(dev,0,0), lddc );
}
}
if(ngpu>1){
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_int_t nbblk = magma_ceildiv((m+blockoffset), nb);
magma_int_t nbblkrow = nbblk-1;
magma_int_t devperm = (dev-stdev+ngpu)%ngpu;
magma_int_t myblk = (nbblkrow/ngpu) + (nbblkrow%ngpu > devperm ? 1:0 );
magma_int_t myrowsize = myblk * nb;
if(dev==stdev) {
myrowsize = myrowsize - blockoffset;
}
//printf("blockoffset %d nbblkrow %d devperm %d DEV %d RECEIVING myblk %d myrowsize %d\n", blockoffset, nbblkrow, devperm, dev, myblk, myrowsize);
if(myrowsize>0){
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
magma_queue_wait_event(streams[ dev ][ 0 ], redevents[dev][1]);
//magma_queue_sync( streams[ dev ][ 1 ] );
// for each dev add the computed ROW block each on its placment with dC
for( magma_int_t blki = 0; blki < myblk; ++blki){
magma_int_t gbblki = (blki*ngpu + devperm)*nb - blockoffset;
magma_int_t lcblki = blki*nb;
magma_int_t ib = nb;// min(nb, m-gbblki);
if(dev==stdev){
lcblki = blki*nb-blockoffset;
if(blki==0){
gbblki = 0;
lcblki = 0;
ib = nb-blockoffset;
}
}
magmablas_sgeadd(ib, n, c_one,
&dwork[dev][lcblki], lddwork,
&dC[dev][gbblki] , lddc );
}
magma_event_record(redevents[dev][0], streams[dev][0]);
}
}
}
// ===========================================================
// COMMUNICATION ALL_REDUCE_SUM
// ===========================================================
if(ngpu==1){
return;
}
// INITIALIZE COMM
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
masterdev = -1;
gnode[cmplxid][MagmaMaxGPUs+1] = -1;
myngpu = gnode[cmplxid][MagmaMaxGPUs];
for( magma_int_t idev = 0; idev < myngpu; ++idev ) {
dev = gnode[cmplxid][idev];
devperm = (dev-stdev+ngpu)%ngpu;
myblk = (nbblk/ngpu) + (nbblk%ngpu > devperm ? 1:0 );
mycolsize = myblk*nb;
myblkoffst = nb*((nbblkoffst/ngpu)+(nbblkoffst%ngpu > dev?1:0));
if(dev==stdev){
mycolsize -= blockoffset;
myblkoffst += blockoffset; // local index in parent matrix
}
if((devperm==devlstblk)&&(lstblksiz>0)){
mycolsize -= (nb-(remm%nb));
}
mycolsize = min(mycolsize, m);
if(mycolsize>0){
gpuisactive[dev] = mycolsize;
if(masterdev==-1) {
masterdev = dev;
nbcmplxactive = nbcmplxactive +1;
cmplxisactive[cmplxid] = 1;
gnode[cmplxid][MagmaMaxGPUs+1] = masterdev;
}
}
}
}
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_device_sync();
}
*/
//*******************************
// each GPU send its result
// to its master. The master make
// the addition and then send to
// to the masters of other real
// and receive from the masters of
// other real make the addition
// and broadcast locally the final
// result.
//*******************************
//printf("=======================================================================\n");
//printf(" sending to my master \n");
//printf("=======================================================================\n");
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//check if real is active
if(masterdev!=-1){
for( magma_int_t idev = 0; idev < myngpu; ++idev ) {
dev = gnode[cmplxid][idev];
mycolsize = gpuisactive[dev];
if(mycolsize>0){
// I am an active GPU. if I am not the master, then send my result to my master.
// store result on dwork[masterdev][dev*maxgsize]
if(dev!=masterdev){
magma_setdevice( dev );
//printf(" GPU %d sending to my master %d\n", dev, masterdev);
// wait the geadd of my ROW and COL GEMM is done
magma_queue_wait_event(streams[ dev ][ 0 ], redevents[dev][0]);
// sending to the master of my real
magma_scopymatrix_async(
m, n,
&dC[dev][0], lddc,
&dwork2[masterdev][maxgsize*dev], m, streams[dev][0] );
magma_event_record(redevents[dev][masterdev], streams[dev][0]);
} // end I am not the masterdev
}// end if mycolsize>0
}// for idev
}// end of if masterdev!=-1 maening real is active
}// for cmplxid
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_device_sync();
}
*/
//printf("=======================================================================\n");
//printf(" each master do addition of local result and broadcast to other masters \n");
//printf("=======================================================================\n");
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//check if real is active
if(masterdev!=-1){
magma_setdevice( masterdev );
// addition is done on stream 0 sequentially
magmablasSetKernelStream( streams[ masterdev ][ 0 ] );
// wait the geadd of my ROW and COL GEMM is done
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[masterdev][0]);
// ========================================
// local addition
// ========================================
for( magma_int_t l = 0; l < myngpu; ++l ) {
lcdev = gnode[cmplxid][l];
lccolsize = gpuisactive[lcdev];
if((lcdev!=masterdev)&&(lccolsize>0)){
//printf(" master %d receiving from %d and adding \n", masterdev, lcdev);
// this is an active GPU of my real.
// wait I received what he send it to me and then do addition.
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[lcdev][masterdev]);
magmablas_sgeadd(m, n, c_one,
&dwork2[masterdev][maxgsize*lcdev], m,
&dC[masterdev][0] , lddc );
}
}// for l=1:myngpu
// because addition is done sequentially on stream 0,
// I have to record this to be able to synch using it
magma_event_record(redevents[masterdev][masterdev], streams[masterdev][0]);
// ========================================
//
// ========================================
// send to other masters
// ========================================
for( magma_int_t k = 0; k < nbcmplx; ++k ) {
if(k!=cmplxid){
gmaster = gnode[k][MagmaMaxGPUs+1];
if(gmaster!=-1){ //real is active
//Master has to wait until finish the local addition then send using gmaster stream.
//use stream 0 to make it sequential or stream gmaster to make it parallel.
//Now both re the same.
//printf(" master %d from cmplx %d sending to other master %d on cmplx %d \n", masterdev, cmplxid, gmaster, k);
magma_queue_wait_event(streams[ masterdev ][ gmaster ], redevents[masterdev][masterdev]);
magma_scopymatrix_async(
m, n,
&dC[masterdev][0], lddc,
&dwork2[gmaster][maxgsize*masterdev], m, streams[masterdev][gmaster] );
magma_event_record(redevents[masterdev][gmaster], streams[masterdev][gmaster]);
magma_event_record(redevents[masterdev][masterdev], streams[masterdev][gmaster]);
} // end of gmaster!=-1
} // end of k!=cmplxid
}// for k = 0: nbcmplx
// ========================================
}// end of if masterdev!=-1 maening real is active
}// for cmplxid
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_device_sync();
}
*/
//printf("=======================================================================\n");
//printf(" each master wait receiving other masters results, do the addition and broadcast locally \n");
//printf("=======================================================================\n");
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//check if real is active
if(masterdev!=-1){
magma_setdevice( masterdev );
// addition is done on stream 0 sequentially
magmablasSetKernelStream( streams[ masterdev ][ 0 ] );
// master has to wait until finishing all the send to other masters.
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[masterdev][masterdev]);
// ========================================
// addition of results from other masters
// ========================================
for( magma_int_t k = 0; k < nbcmplx; ++k ) {
if(k!=cmplxid){
gmaster = gnode[k][MagmaMaxGPUs+1];
if(gmaster!=-1){ //real is active
//Master has to wait until receiving from gmaster, then do addition using stream 0
//printf(" master %d from cmplx %d receiving from other master %d on cmplx %d and adding \n", masterdev, cmplxid, gmaster, k);
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[gmaster][masterdev]);
magmablas_sgeadd(m, n, c_one,
&dwork2[masterdev][maxgsize*gmaster], m,
&dC[masterdev][0] , lddc );
} // end of gmaster!=-1
} // end of k!=cmplxid
}// for k = 0: nbcmplx
// because addition is done sequentially on stream 0,
// I have to record this to be able to synch using it
magma_event_record(redevents[masterdev][masterdev], streams[masterdev][0]);
// ========================================
// ========================================
// local broadcast of final results
// ========================================
for( magma_int_t l = 0; l < myngpu; ++l ) {
lcdev = gnode[cmplxid][l];
lccolsize = gpuisactive[lcdev];
if((lcdev!=masterdev)&&(lccolsize>0)){
// this is an active GPU of my real.
// wait the previous addition is done maening stream 0 is finished and broadcast sequentially for now.
// to make it parallel put stream lcdev instead of stream 0
//printf(" master %d broadcasting local to %d \n", masterdev, lcdev);
magma_queue_wait_event(streams[ masterdev ][ 0 ], redevents[masterdev][masterdev]);
magma_scopymatrix_async(
m, n,
&dC[masterdev][0], lddc,
&dC[lcdev][0], lddc, streams[masterdev][0] );
magma_event_record(redevents[masterdev][lcdev], streams[masterdev][0]);
}
}// for l=1:myngpu
// ========================================
}// end of if masterdev!=-1 maening real is active
}// for cmplxid
/*
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_device_sync();
}
*/
for( magma_int_t cmplxid = 0; cmplxid < nbcmplx; ++cmplxid ) {
myngpu = gnode[cmplxid][MagmaMaxGPUs];
masterdev = gnode[cmplxid][MagmaMaxGPUs+1];
//check if real is active
if(masterdev!=-1){
for( magma_int_t l = 0; l < myngpu; ++l ) {
lcdev = gnode[cmplxid][l];
lccolsize = gpuisactive[lcdev];
if(lccolsize>0){
magma_setdevice( lcdev );
magma_queue_wait_event(streams[ lcdev ][ 0 ], redevents[lcdev][0]);
magma_queue_wait_event(streams[ lcdev ][ 0 ], redevents[masterdev][lcdev]);
}
}// for l=1:myngpu
}// end of if masterdev!=-1 maening real is active
}// for cmplxid
//printf("****************************************************\n");
//printf(" finish ssymm \n");
//printf("****************************************************\n");
magma_setdevice( cdev );
magmablasSetKernelStream( cstream );
}
extern "C" void
magma_zherk_mgpu(
magma_int_t ngpu,
magma_uplo_t uplo, magma_trans_t trans, magma_int_t nb, magma_int_t n, magma_int_t k,
double alpha,
magmaDoubleComplex_ptr dB[], magma_int_t lddb, magma_int_t b_offset,
double beta,
magmaDoubleComplex_ptr dC[], magma_int_t lddc, magma_int_t c_offset,
magma_int_t nqueue, magma_queue_t queues[][10])
{
#define dB(id, i, j) (dB[(id)]+(j)*lddb + (i)+b_offset)
#define dC(id, i, j) (dC[(id)]+(j)*lddc + (i))
#define STREAM_ID(i) (nqueue > 1 ? 1+((i)/nb)%(nqueue-1) : 0)
magma_int_t i, id, ib, ii, kk, n1;
magmaDoubleComplex z_alpha = MAGMA_Z_MAKE(alpha,0.0);
magmaDoubleComplex z_beta = MAGMA_Z_MAKE(beta, 0.0);
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
/* diagonal update */
for( i=0; i < n; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = STREAM_ID( i+c_offset );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
/* zher2k on diagonal block */
magma_setdevice(id);
magmablasSetKernelStream( queues[id][kk] );
trace_gpu_start( id, kk, "syr2k", "syr2k" );
magma_zherk(uplo, trans, ib, k,
alpha, dB(id, i, 0 ), lddb,
beta, dC(id, i+c_offset, ii), lddc);
trace_gpu_end( id, kk );
}
/* off-diagonal update */
if (uplo == MagmaUpper) {
for( i=nb; i < n; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = STREAM_ID( i+c_offset );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
magma_setdevice(id);
magmablasSetKernelStream( queues[id][kk] );
magma_zgemm(MagmaNoTrans, MagmaConjTrans, i, ib, k,
z_alpha, dB(id, 0, 0 ), lddb,
dB(id, i, 0 ), lddb,
z_beta, dC(id, 0, ii), lddc);
}
}
else {
for( i=0; i < n-nb; i += nb ) {
id = ((i+c_offset)/nb)%ngpu;
kk = STREAM_ID( i+c_offset );
ib = min(nb, n-i);
ii = nb*((i+c_offset)/(nb*ngpu));
n1 = n-i-ib;
/* zgemm on off-diagonal blocks */
magma_setdevice(id);
magmablasSetKernelStream( queues[id][kk] );
trace_gpu_start( id, kk, "gemm_up", "gemm_up" );
magma_zgemm(MagmaNoTrans, MagmaConjTrans, n1, ib, k,
z_alpha, dB(id, i+ib, 0 ), lddb,
dB(id, i, 0 ), lddb,
z_beta, dC(id, i+c_offset+ib, ii), lddc);
trace_gpu_end( id, kk );
}
}
// TODO why not sync?
//for( id=0; id < ngpu; id++ ) {
// magma_setdevice(id);
// //for( kk=0; kk < nqueue; kk++ )
// // magma_queue_sync( queues[id][kk] );
//}
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
}
