// use_cuda_time = 1: use cudaEventElapsedTime()
// or use getSystemTime()
void test_2gpu(float *d_send_data, float *d_recv_data, int size, int id0, int id1, bool use_cuda_time)
{
if(use_cuda_time) {
cudaEvent_t start_event, stop_event;
float time_memcpy;
// version I
//cudaEventCreate(&start_event);
//cudaEventCreate(&stop_event);
//cudaEventRecord(start_event, 0);
// version II
int eventflags = cudaEventBlockingSync;
cudaEventCreateWithFlags(&start_event, eventflags);
cudaEventCreateWithFlags(&stop_event, eventflags);
cudaEventRecord(start_event, 0);
for(int i=0; i<CNT; i++) {
cudaMemcpy(d_recv_data, d_send_data, size*sizeof(float), cudaMemcpyDeviceToDevice);
}
std::cout << "hello, use_cuda_time" << std::endl;
cudaEventRecord(stop_event, 0);
cudaEventSynchronize(stop_event);
cudaEventElapsedTime(&time_memcpy, start_event, stop_event); // ms
std::cout << "Time is " << time_memcpy/1000. << "s" << std::endl;
std::cout << "GPU" << id0 << " ---> GPU" << id1 << " :" <<
WIDTH*HEIGHT*sizeof(float)*CNT*1000./(1024*1024*time_memcpy) << "MB/s" << std::endl;
cudaEventDestroy(start_event);
cudaEventDestroy(stop_event);
} else {
//cudaEvent_t start_event;
//cudaEventCreate(&start_event);
long long start = getSystemTime();
for(int i=0; i<CNT; i++) {
cudaMemcpy(d_recv_data, d_send_data, size*sizeof(float), cudaMemcpyDeviceToDevice);
//cudaMemcpyPeer(d_recv_data, id1, d_send_data, id0, size*sizeof(float));
}
//cudaEventRecord(start_event, 0);
//cudaEventSynchronize(start_event);
long long end = getSystemTime();
std::cout << "Time is " << (end-start)/1000. << "s" << std::endl;
std::cout << "GPU" << id0 << " ---> GPU" << id1 << " :" <<
WIDTH*HEIGHT*sizeof(float)*CNT*1000./(1024*1024*(end - start+1)) << "MB/s" << std::endl;
} //WIDTH*HEIGHT*4.*CNT/(1000*(end - start)) << "Mb/s" << std::endl;
}
void blasx_resource_init(int GPUs, cublasHandle_t* handles, cudaStream_t* streams, cudaEvent_t* events, void** C_dev, int floatType_id) {
if(floatType_id == 0) C_dev = (float**) C_dev;
else if(floatType_id == 1) C_dev = (double**) C_dev;
else C_dev = (cuDoubleComplex**) C_dev;
int GPU_id = 0;
for (GPU_id = 0; GPU_id < GPUs; GPU_id++) {
assert( cudaSetDevice(GPU_id) == cudaSuccess );
//create handles
assert( cublasCreate(&handles[GPU_id]) == CUBLAS_STATUS_SUCCESS);
//create streams and event
int i = 0;
for (i = 0 ; i < STREAMNUM; i++) {
assert( cudaStreamCreate(&streams[i+GPU_id*STREAMNUM]) == cudaSuccess );
assert( cudaEventCreateWithFlags(&events[i+GPU_id*STREAMNUM], cudaEventDisableTiming) == cudaSuccess );
}
//create C_dev
for (i = 0; i < STREAMNUM*2; i++) {
if (floatType_id == 0) {
assert( cudaMalloc((void**)&C_dev[i+GPU_id*STREAMNUM*2], sizeof(float)*BLOCKSIZE_SGEMM*BLOCKSIZE_SGEMM) == cudaSuccess );
}else if (floatType_id == 1) {
assert( cudaMalloc((void**)&C_dev[i+GPU_id*STREAMNUM*2], sizeof(double)*BLOCKSIZE_DGEMM*BLOCKSIZE_DGEMM) == cudaSuccess );
} else {
assert( cudaMalloc((void**)&C_dev[i+GPU_id*STREAMNUM*2], sizeof(cuDoubleComplex)*BLOCKSIZE_ZGEMM*BLOCKSIZE_ZGEMM) == cudaSuccess );
}
}
}
}
/**
* create valid object
*
* - internal memory is allocated
* - event must be destroyed with @see destroy
*/
static CudaEvent create()
{
CudaEvent ev;
ev.isValid = true;
CUDA_CHECK(cudaEventCreateWithFlags(&(ev.event), cudaEventDisableTiming));
return ev;
}
GPUDataTransferer<ElemType>::GPUDataTransferer(int deviceId, bool useConcurrentStreams)
: m_deviceId(deviceId)
{
PrepareDevice(m_deviceId);
// events
// Note: Do NOT use cudaEventBlockingSync (which supposedly yields the process)--it will totally break cudaEventSynchronize(), causing it to take 50 or 100 ms randomly.
cudaEventCreateWithFlags(&m_fetchCompleteEvent, cudaEventDisableTiming) || "cudaEventCreateWithFlags failed";
cudaEventCreateWithFlags(&m_assignCompleteEvent, cudaEventDisableTiming) || "cudaEventCreateWithFlags failed";
#pragma warning(disable : 4127)
if (useConcurrentStreams && (m_fetchStream == NULL))
{
cudaStreamCreateWithFlags(&m_fetchStream, cudaStreamNonBlocking) || "cudaStreamCreateWithFlags failed";
cudaStreamCreateWithFlags(&m_assignStream, cudaStreamNonBlocking) || "cudaStreamCreateWithFlags failed";
}
}
AsyncCopier::Event::Event(int d)
: device(d),
refCount(0) {
event.emplace();
checkCudaError(
cudaEventCreateWithFlags(
get_pointer(event), cudaEventDisableTiming | cudaEventBlockingSync),
"cudaEventCreateWithFlags");
}
__host__ __device__
future(cudaStream_t s, bool owns_stream)
: m_stream(s),m_owns_stream(owns_stream)
{
#if __BULK_HAS_CUDART__
bulk::detail::throw_on_error(cudaEventCreateWithFlags(&m_event, create_flags), "cudaEventCreateWithFlags in future ctor");
bulk::detail::throw_on_error(cudaEventRecord(m_event, m_stream), "cudaEventRecord in future ctor");
#endif
} // end future()
/**
* Adds count new cuda events to the pool.
* @param count number of cuda events to add
*/
void addEvents(size_t count)
{
for (size_t i = 0; i < count; i++)
{
cudaEvent_t event;
CUDA_CHECK(cudaEventCreateWithFlags(&event,cudaEventDisableTiming));
events.push_back(event);
}
}
/*
* Initializes the array of CUDA Events for timing.
*
* Returns EXIT_SUCCESS OR EXIT_FAILURE.
*/
int init_timing_events( void )
{
#if NMFGPU_PROFILING_TRANSF || NMFGPU_PROFILING_KERNELS
#if NMFGPU_VERBOSE_2
print_message( verb_shown_by_all, "Initializing array of CUDA Events for timing (number of events: %" PRI_IDX ")...\n",
NUM_TIMING_EVENTS );
#endif
cudaError_t cuda_status = cudaSuccess;
// ----------------------------
// Start timer
cuda_status = cudaEventCreateWithFlags( &timing_events[ START_EVENT ], cudaEventBlockingSync );
if ( cuda_status != cudaSuccess ) {
print_error( sys_error_shown_by_all, "Error creating CUDA event for timing (timer start): %s\n",
cudaGetErrorString(cuda_status) );
return EXIT_FAILURE;
}
// Stop timer
cuda_status = cudaEventCreateWithFlags( &timing_events[ STOP_EVENT ], cudaEventBlockingSync );
if ( cuda_status != cudaSuccess ) {
print_error( sys_error_shown_by_all, "Error creating CUDA event for timing (timer stop): %s\n",
cudaGetErrorString(cuda_status) );
cudaEventDestroy( timing_events[ START_EVENT ] );
return EXIT_FAILURE;
}
// ----------------------------
#if NMFGPU_VERBOSE_2
print_message( verb_shown_by_all, "Initializing array of CUDA Events for timing (number of events: %"
PRI_IDX ")... Done.\n", NUM_TIMING_EVENTS );
#endif
#endif /* if NMFGPU_PROFILING_TRANSF || NMFGPU_PROFILING_KERNELS */
return EXIT_SUCCESS;
} // init_timing_events
//-----------------------------------------------------------------------------
//! Constructor.
//-----------------------------------------------------------------------------
ALPAKA_FN_HOST EventCudaImpl(
dev::DevCudaRt const & dev,
bool bBusyWait) :
m_Dev(dev),
m_CudaEvent()
{
ALPAKA_DEBUG_MINIMAL_LOG_SCOPE;
// Set the current device.
ALPAKA_CUDA_RT_CHECK(cudaSetDevice(
m_Dev.m_iDevice));
// Create the event on the current device with the specified flags. Valid flags include:
// - cudaEventDefault: Default event creation flag.
// - cudaEventBlockingSync : Specifies that event should use blocking synchronization.
// A host thread that uses cudaEventSynchronize() to wait on an event created with this flag will block until the event actually completes.
// - cudaEventDisableTiming : Specifies that the created event does not need to record timing data.
// Events created with this flag specified and the cudaEventBlockingSync flag not specified will provide the best performance when used with cudaStreamWaitEvent() and cudaEventQuery().
ALPAKA_CUDA_RT_CHECK(cudaEventCreateWithFlags(
&m_CudaEvent,
(bBusyWait ? cudaEventDefault : cudaEventBlockingSync) | cudaEventDisableTiming));
}
int createSingleDeviceEvents(lua_State *L, THCState *state, int arg,
int device, cudaEvent_t* event)
{
/* Push table to top */
lua_pushvalue(L, arg);
/* Record events */
lua_pushnil(L);
int i = 0;
while (lua_next(L, -2)) {
int streamId = (int) lua_tonumber(L, -1);
cudaStream_t streamWaitingOn =
THCState_getDeviceStream(state, device, streamId);
THCudaCheck(cudaEventCreateWithFlags(&event[i], cudaEventDisableTiming));
THCudaCheck(cudaEventRecord(event[i], streamWaitingOn));
lua_pop(L, 1);
i++;
}
/* Pop table from top */
lua_pop(L, 1);
return i;
}
int TEMPLATE2 (CHOLMOD (gpu_init))
(
void *Cwork,
cholmod_factor *L,
cholmod_common *Common,
Int nsuper,
Int n,
Int nls,
cholmod_gpu_pointers *gpu_p
)
{
Int i, k, maxSize ;
cublasStatus_t cublasError ;
cudaError_t cudaErr ;
size_t maxBytesSize, HostPinnedSize ;
feenableexcept (FE_DIVBYZERO | FE_INVALID | FE_OVERFLOW );
maxSize = L->maxcsize;
/* #define PAGE_SIZE (4*1024) */
CHOLMOD_GPU_PRINTF (("gpu_init : %p\n",
(void *) ((size_t) Cwork & ~(4*1024-1)))) ;
/* make sure the assumed buffer sizes are large enough */
if ( (nls+2*n+4)*sizeof(Int) > Common->devBuffSize ) {
ERROR (CHOLMOD_GPU_PROBLEM,"\n\n"
"GPU Memory allocation error. Ls, Map and RelativeMap exceed\n"
"devBuffSize. It is not clear if this is due to insufficient\n"
"device or host memory or both. You can try:\n"
" 1) increasing the amount of GPU memory requested\n"
" 2) reducing CHOLMOD_NUM_HOST_BUFFERS\n"
" 3) using a GPU & host with more memory\n"
"This issue is a known limitation and should be fixed in a \n"
"future release of CHOLMOD.\n") ;
return (0) ;
}
/* divvy up the memory in dev_mempool */
gpu_p->d_Lx[0] = Common->dev_mempool;
gpu_p->d_Lx[1] = Common->dev_mempool + Common->devBuffSize;
gpu_p->d_C = Common->dev_mempool + 2*Common->devBuffSize;
gpu_p->d_A[0] = Common->dev_mempool + 3*Common->devBuffSize;
gpu_p->d_A[1] = Common->dev_mempool + 4*Common->devBuffSize;
gpu_p->d_Ls = Common->dev_mempool + 5*Common->devBuffSize;
gpu_p->d_Map = gpu_p->d_Ls + (nls+1)*sizeof(Int) ;
gpu_p->d_RelativeMap = gpu_p->d_Map + (n+1)*sizeof(Int) ;
/* Copy all of the Ls and Lpi data to the device. If any supernodes are
* to be computed on the device then this will be needed, so might as
* well do it now. */
cudaErr = cudaMemcpy ( gpu_p->d_Ls, L->s, nls*sizeof(Int),
cudaMemcpyHostToDevice );
CHOLMOD_HANDLE_CUDA_ERROR(cudaErr,"cudaMemcpy(d_Ls)");
if (!(Common->gpuStream[0])) {
/* ------------------------------------------------------------------ */
/* create each CUDA stream */
/* ------------------------------------------------------------------ */
for ( i=0; i<CHOLMOD_HOST_SUPERNODE_BUFFERS; i++ ) {
cudaErr = cudaStreamCreate ( &(Common->gpuStream[i]) );
if (cudaErr != cudaSuccess) {
ERROR (CHOLMOD_GPU_PROBLEM, "CUDA stream") ;
return (0) ;
}
}
/* ------------------------------------------------------------------ */
/* create each CUDA event */
/* ------------------------------------------------------------------ */
for (i = 0 ; i < 3 ; i++) {
cudaErr = cudaEventCreateWithFlags
(&(Common->cublasEventPotrf [i]), cudaEventDisableTiming) ;
if (cudaErr != cudaSuccess) {
ERROR (CHOLMOD_GPU_PROBLEM, "CUDA event") ;
return (0) ;
}
}
for (i = 0 ; i < CHOLMOD_HOST_SUPERNODE_BUFFERS ; i++) {
cudaErr = cudaEventCreateWithFlags
(&(Common->updateCBuffersFree[i]), cudaEventDisableTiming) ;
if (cudaErr != cudaSuccess) {
ERROR (CHOLMOD_GPU_PROBLEM, "CUDA event") ;
return (0) ;
}
}
cudaErr = cudaEventCreateWithFlags ( &(Common->updateCKernelsComplete),
cudaEventDisableTiming );
if (cudaErr != cudaSuccess) {
ERROR (CHOLMOD_GPU_PROBLEM, "CUDA updateCKernelsComplete event") ;
return (0) ;
}
}
gpu_p->h_Lx[0] = (double*)(Common->host_pinned_mempool);
for ( k=1; k<CHOLMOD_HOST_SUPERNODE_BUFFERS; k++ ) {
gpu_p->h_Lx[k] = (double*)((char *)(Common->host_pinned_mempool) +
k*Common->devBuffSize);
}
return (1); /* initialization successfull, useGPU = 1 */
}
/**
Purpose
-------
DSYTRD_HE2HB reduces a real symmetric matrix A to real symmetric
band-diagonal form T by an orthogonal similarity transformation:
Q**H * A * Q = T.
This version stores the triangular matrices T used in the accumulated
Householder transformations (I - V T V').
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. 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.
On exit, if UPLO = MagmaUpper, the Upper band-diagonal of A is
overwritten by the corresponding elements of the
band-diagonal matrix T, and the elements above the band
diagonal, with the array TAU, represent the orthogonal
matrix Q as a product of elementary reflectors; if UPLO
= MagmaLower, the the Lower band-diagonal of A is overwritten by
the corresponding elements of the band-diagonal
matrix T, and the elements below the band-diagonal, with
the array TAU, represent the orthogonal matrix Q as a product
of elementary reflectors. See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
tau DOUBLE_PRECISION array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) DOUBLE_PRECISION array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= 1.
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
dT DOUBLE_PRECISION array on the GPU, dimension N*NB,
where NB is the optimal blocksize.
On exit dT holds the upper triangular matrices T from the
accumulated Householder transformations (I - V T V') used
in the factorization. The nb x nb matrices T are ordered
consecutively in memory one after another.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n-1) . . . H(2) H(1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(i+1:n) = 0 and v(i) = 1; v(1:i-1) is stored on exit in
A(1:i-1,i+1), and tau in TAU(i).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(n-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i) = 0 and v(i+1) = 1; v(i+2:n) is stored on exit in A(i+2:n,i),
and tau in TAU(i).
The contents of A on exit are illustrated by the following examples
with n = 5:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( d e v2 v3 v4 ) ( d )
( d e v3 v4 ) ( e d )
( d e v4 ) ( v1 e d )
( d e ) ( v1 v2 e d )
( d ) ( v1 v2 v3 e d )
where d and e denote diagonal and off-diagonal elements of T, and vi
denotes an element of the vector defining H(i).
@ingroup magma_dsyev_2stage
********************************************************************/
extern "C" magma_int_t
magma_dsytrd_sy2sb_mgpu(
magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
double *A, magma_int_t lda,
double *tau,
double *work, magma_int_t lwork,
magmaDouble_ptr dAmgpu[], magma_int_t ldda,
magmaDouble_ptr dTmgpu[], magma_int_t lddt,
magma_int_t ngpu, magma_int_t distblk,
magma_queue_t queues[][20], magma_int_t nqueue,
magma_int_t *info)
{
#define A(a_1,a_2) ( A + ((a_2)-1)*( lda) + (a_1)-1)
#define tau_ref(a_1) (tau + (a_1)-1)
#define dT(a_0, a_1, a_2) (dTmgpu[a_0] + ((a_2)-1)*(lddt) + (a_1)-1)
#define dA(a_0, a_1, a_2) (dAmgpu[a_0] + ((a_2)-1)*(ldda) + (a_1)-1)
double c_neg_one = MAGMA_D_NEG_ONE;
double c_neg_half = MAGMA_D_NEG_HALF;
double c_one = MAGMA_D_ONE;
double c_zero = MAGMA_D_ZERO;
double d_one = MAGMA_D_ONE;
magma_int_t pm, pn, indi, indj, pk;
magma_int_t pm_old=0, pn_old=0, indi_old=0, flipV=-1;
magma_int_t iblock, idev, di;
int i;
int lwkopt;
int lquery;
assert (nqueue >= 3);
assert (nqueue >= (ngpu+1));
*info = 0;
int upper = (uplo == MagmaUpper);
lquery = (lwork == -1);
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
} else if (lwork < 1 && ! lquery) {
*info = -9;
}
/* Determine the block size. */
lwkopt = n * nb;
if (*info == 0) {
work[0] = MAGMA_D_MAKE( lwkopt, 0 );
}
if (*info != 0)
return *info;
else if (lquery)
return *info;
/* Quick return if possible */
if (n == 0) {
work[0] = c_one;
return *info;
}
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
// limit to 16 threads
magma_int_t orig_threads = magma_get_lapack_numthreads();
magma_set_lapack_numthreads( min(orig_threads,16) );
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2];
magma_int_t nbcmplx=0;
magma_buildconnection_mgpu(gnode, &nbcmplx, ngpu);
#ifdef ENABLE_DEBUG
printf(" Initializing communication pattern.... GPU-ncmplx %d\n\n", nbcmplx);
#endif
double *dspace[MagmaMaxGPUs];
double *dwork[MagmaMaxGPUs], *dworkbis[MagmaMaxGPUs];
double *dvall[MagmaMaxGPUs], *dv[MagmaMaxGPUs], *dw[MagmaMaxGPUs];
double *workngpu[MagmaMaxGPUs+1];
magma_event_t redevents[MagmaMaxGPUs][MagmaMaxGPUs*MagmaMaxGPUs+10];
magma_int_t nbevents = MagmaMaxGPUs*MagmaMaxGPUs;
magma_int_t lddv = ldda;
magma_int_t lddw = lddv;
magma_int_t dwrk2siz = ldda*nb*(ngpu+1);
magma_int_t worksiz = n*nb;
magma_int_t devworksiz = 2*nb*lddv + nb*lddw + nb*ldda + dwrk2siz; // 2*dv(dv0+dv1) + dw + dwork +dworkbis
// local allocation and stream creation
// TODO check malloc
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_dmalloc( &dspace[dev], devworksiz );
magma_dmalloc_pinned ( &workngpu[dev], worksiz);
dvall[dev] = dspace[dev];
dw[dev] = dvall[dev] + 2*nb*lddv;
dwork[dev] = dw[dev] + nb*lddw;
dworkbis[dev] = dwork[dev] + nb*ldda;
magmablasSetKernelStream( queues[ dev ][ 0 ] );
for( magma_int_t i = 0; i < nbevents; ++i ) {
cudaEventCreateWithFlags(&redevents[dev][i],cudaEventDisableTiming);
}
}
magma_dmalloc_pinned ( &workngpu[ngpu], worksiz);
double *worktest = NULL;
//magma_dmalloc_cpu( &worktest, n*nb ); // not used
// ======================
double *hT = work + lwork - nb*nb;
lwork -= nb*nb;
memset( hT, 0, nb*nb*sizeof(double));
if (upper) {
printf("DSYTRD_HE2HB is not yet implemented for upper matrix storage. Exit.\n");
exit(1);
} else {
/* Reduce the lower triangle of A */
for (i = 1; i <= n-nb; i += nb) {
indi = i+nb;
indj = i;
pm = n - i - nb + 1;
//pn = min(i+nb-1, n-nb) -i + 1;
pn = nb;
/* Get the current panel (no need for the 1st iteration) */
if (i > 1 ) {
// dpanel_to_q copy the upper oof diagonal part of
// the matrix to work to be restored later. acctually
// the zero's and one's putted are not used this is only
// because we don't have a function that copy only the
// upper part of A to be restored after copying the
// lookahead panel that has been computted from GPU to CPU.
dpanel_to_q(MagmaUpper, pn-1, A(i, i+1), lda, work);
// find the device who own the panel then send it to the CPU.
// below a -1 was added and then a -1 was done on di because of the fortran indexing
iblock = ((i-1) / distblk) / ngpu; // local block id
di = iblock*distblk + (i-1)%distblk; // local index in parent matrix
idev = ((i-1) / distblk) % ngpu; // device with this block
//printf("Receiving panel ofsize %d %d from idev %d A(%d,%d) \n",(pm+pn), pn,idev,i-1,di);
magma_setdevice( idev );
//magma_device_sync();
magma_dgetmatrix_async( (pm+pn), pn,
dA(idev, i, di+1), ldda,
A( i, i), lda, queues[ idev ][ nqueue-1 ] );
//magma_setdevice( 0 );
//printf("updating dsyr2k on A(%d,%d) of size %d %d \n",indi_old+pn_old-1,indi_old+pn_old-1,pm_old-pn_old,pn_old);
// compute DSYR2K_MGPU
magmablas_dsyr2k_mgpu2(
MagmaLower, MagmaNoTrans, pm_old-pn_old, pn_old,
c_neg_one, dv, pm_old, pn_old,
dw, pm_old, pn_old,
d_one, dAmgpu, ldda, indi_old+pn_old-1,
ngpu, distblk, queues, 2 );
//magma_setdevice( 0 );
magma_setdevice( idev );
magma_queue_sync( queues[idev][ nqueue-1 ] );
//magma_setdevice( 0 );
dq_to_panel(MagmaUpper, pn-1, A(i, i+1), lda, work);
}
/* ==========================================================
QR factorization on a panel starting nb off of the diagonal.
Prepare the V and T matrices.
========================================================== */
lapackf77_dgeqrf(&pm, &pn, A(indi, indj), &lda,
tau_ref(i), work, &lwork, info);
/* Form the matrix T */
pk=min(pm,pn);
lapackf77_dlarft( MagmaForwardStr, MagmaColumnwiseStr,
&pm, &pk, A(indi, indj), &lda,
tau_ref(i), hT, &nb);
/* Prepare V - put 0s in the upper triangular part of the panel
(and 1s on the diagonal), temporaly storing the original in work */
dpanel_to_q(MagmaUpper, pk, A(indi, indj), lda, work);
/* Send V and T from the CPU to the GPU */
// To be able to overlap the GET with the DSYR2K
// it should be done on last stream.
// TO Avoid a BUG that is overwriting the old_V
// used atthis moment by dsyr2k with the new_V
// send it now, we decide to have a flipflop
// vector of Vs. if step%2=0 use V[0] else use V[nb*n]
flipV = ((i-1)/nb)%2;
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
dv[dev] = dvall[dev] + flipV*nb*lddv;
}
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
// send V
magma_dsetmatrix_async( pm, pk,
A(indi, indj), lda,
dv[dev], pm, queues[dev][nqueue-1] );
// Send the triangular factor T to the GPU
magma_dsetmatrix_async( pk, pk,
hT, nb,
dT(dev, 1, i), lddt, queues[dev][nqueue-1] );
}
/* ==========================================================
Compute W:
1. X = A (V T)
2. W = X - 0.5* V * (T' * (V' * X))
========================================================== */
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
// dwork = V T
magma_setdevice( dev );
magmablasSetKernelStream( queues[ dev ][ nqueue-1 ] );
magma_queue_sync( queues[dev][nqueue-1] );
magma_dgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_one, dv[dev], pm,
dT(dev, 1, i), lddt,
c_zero, dwork[dev], pm);
}
// ===============================================
// SYNC TO BE SURE THAT BOTH V AND T WERE
// RECEIVED AND VT IS COMPUTED and SYR2K is done
// ===============================================
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
for( magma_int_t s = 0; s < nqueue; ++s )
magma_queue_sync( queues[dev][s] );
}
// compute DSYMM_MGPU
// The broadcast of the result done inside this function
// should be done in stream [0] because i am assuming this
// for the GEMMs below otherwise I have to SYNC over the
// Broadcasting stream.
if (ngpu == 1) {
magmablasSetKernelStream( queues[ 0 ][ 0 ] );
magma_dsymm(MagmaLeft, uplo, pm, pk,
c_one, dAmgpu[0]+(indi-1)*ldda+(indi-1), ldda,
dwork[0], pm,
c_zero, dw[0], pm);
} else {
magmablas_dsymm_mgpu_com(
MagmaLeft, uplo, pm, pk,
c_one, dAmgpu, ldda, indi-1,
dwork, pm,
c_zero, dw, pm, dworkbis, dwrk2siz, worktest, pm, workngpu, worksiz,
ngpu, distblk, queues, nqueue-1, redevents, nbevents, gnode, nbcmplx);
}
/* dwork = V*T already ==> dwork' = T'*V'
* compute T'*V'*X ==> dwork'*W ==>
* dwork + pm*nb = ((T' * V') * X) = dwork' * X = dwork' * W */
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
// Here we have to wait until the broadcast of DSYMM has been done.
// Note that the broadcast should be done on stream[0] so in a way
// we can continue here on the same stream and avoid a sync
magma_setdevice( dev );
magmablasSetKernelStream( queues[ dev ][ 0 ] );
// magma_queue_sync( queues[dev][0] );
magma_dgemm(MagmaConjTrans, MagmaNoTrans, pk, pk, pm,
c_one, dwork[dev], pm,
dw[dev], pm,
c_zero, dworkbis[dev], nb);
/* W = X - 0.5 * V * T'*V'*X
* = X - 0.5 * V * (dwork + pm*nb) = W - 0.5 * V * (dwork + pm*nb) */
magma_dgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_neg_half, dv[dev], pm,
dworkbis[dev], nb,
c_one, dw[dev], pm);
}
/* restore the panel it is put here to overlap with the previous GEMM*/
dq_to_panel(MagmaUpper, pk, A(indi, indj), lda, work);
// ===============================================
// SYNC TO BE SURE THAT BOTH V AND W ARE DONE
// ===============================================
// Synchronise to be sure that W has been computed
// because next DSYR2K use streaming and may happen
// that lunch a gemm on stream 2 while stream 0
// which compute those 2 GEMM above has not been
// computed and also used for the same reason in
// the panel update below and also for the last HER2K
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_queue_sync( queues[dev][0] );
}
/* ==========================================================
Update the unreduced submatrix A(i+ib:n,i+ib:n), using
an update of the form: A := A - V*W' - W*V'
========================================================== */
if (i + nb <= n-nb) {
/* There would be next iteration;
do lookahead - update the next panel */
// below a -1 was added and then a -1 was done on di because of the fortran indexing
iblock = ((indi-1) / distblk) / ngpu; // local block id
di = iblock*distblk + (indi-1)%distblk; // local index in parent matrix
idev = ((indi-1) / distblk) % ngpu; // device with this block
magma_setdevice( idev );
magmablasSetKernelStream( queues[ idev ][ nqueue-1 ] );
//magma_queue_sync( queues[idev][0] ); removed because the sync has been done in the loop above
magma_dgemm(MagmaNoTrans, MagmaConjTrans, pm, pn, pn, c_neg_one,
dv[idev], pm,
dw[idev], pm, c_one,
dA(idev, indi, di+1), ldda);
magma_dgemm(MagmaNoTrans, MagmaConjTrans, pm, pn, pn, c_neg_one,
dw[idev], pm,
dv[idev], pm, c_one,
dA(idev, indi, di+1), ldda);
//printf("updating next panel distblk %d idev %d on A(%d,%d) of size %d %d %d \n",distblk,idev,indi-1,di,pm,pn,pn);
}
else {
/* no look-ahead as this is last iteration */
// below a -1 was added and then a -1 was done on di because of the fortran indexing
iblock = ((indi-1) / distblk) / ngpu; // local block id
di = iblock*distblk + (indi-1)%distblk; // local index in parent matrix
idev = ((indi-1) / distblk) % ngpu; // device with this block
magma_setdevice( idev );
magmablasSetKernelStream( queues[ idev ][ 0 ] );
//printf("LAST DSYR2K idev %d on A(%d,%d) of size %d \n",idev, indi-1,di,pk);
magma_dsyr2k(MagmaLower, MagmaNoTrans, pk, pk, c_neg_one,
dv[idev], pm,
dw[idev], pm, d_one,
dA(idev, indi, di+1), ldda);
/* Send the last block to the CPU */
dpanel_to_q(MagmaUpper, pk-1, A(n-pk+1, n-pk+2), lda, work);
magma_dgetmatrix( pk, pk,
dA(idev, indi, di+1), ldda,
A(n-pk+1, n-pk+1), lda );
dq_to_panel(MagmaUpper, pk-1, A(n-pk+1, n-pk+2), lda, work);
}
indi_old = indi;
//indj_old = indj;
pm_old = pm;
pn_old = pn;
} // end loop for (i)
}// end of LOWER
//magma_setdevice( 0 );
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_free( dspace[dev]);
magma_free_pinned(workngpu[dev]);
for( magma_int_t e = 0; e < nbevents; ++e ) {
magma_event_destroy( redevents[dev][e] );
}
}
magma_free_pinned(workngpu[ngpu]);
magma_free_cpu(worktest);
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
magma_set_lapack_numthreads( orig_threads );
work[0] = MAGMA_D_MAKE( lwkopt, 0 );
return *info;
} /* magma_dsytrd_sy2sb_mgpu */
extern "C" magma_int_t
magma_chetrd_he2hb_mgpu( char uplo, magma_int_t n, magma_int_t nb,
magmaFloatComplex *a, magma_int_t lda,
magmaFloatComplex *tau,
magmaFloatComplex *work, magma_int_t lwork,
magmaFloatComplex *dAmgpu[], magma_int_t ldda,
magmaFloatComplex *dTmgpu[], magma_int_t lddt,
magma_int_t ngpu, magma_int_t distblk,
magma_queue_t streams[][20], magma_int_t nstream,
magma_int_t threads, magma_int_t *info)
{
/* -- MAGMA (version 1.4.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
August 2013
Purpose
=======
CHETRD_HE2HB reduces a complex Hermitian matrix A to real symmetric
band-diagonal form T by an orthogonal similarity transformation:
Q**H * A * Q = T.
This version stores the triangular matrices T used in the accumulated
Householder transformations (I - V T V').
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
A (input/output) COMPLEX 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 UPLO = 'U', the Upper band-diagonal of A is
overwritten by the corresponding elements of the
band-diagonal matrix T, and the elements above the band
diagonal, with the array TAU, represent the orthogonal
matrix Q as a product of elementary reflectors; if UPLO
= 'L', the the Lower band-diagonal of A is overwritten by
the corresponding elements of the band-diagonal
matrix T, and the elements below the band-diagonal, with
the array TAU, represent the orthogonal matrix Q as a product
of elementary reflectors. See Further Details.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
TAU (output) COMPLEX array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details).
WORK (workspace/output) COMPLEX array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The dimension of the array WORK. LWORK >= 1.
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
dT (output) COMPLEX array on the GPU, dimension N*NB,
where NB is the optimal blocksize.
On exit dT holds the upper triangular matrices T from the
accumulated Householder transformations (I - V T V') used
in the factorization. The nb x nb matrices T are ordered
consecutively in memory one after another.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
Further Details
===============
If UPLO = 'U', the matrix Q is represented as a product of elementary
reflectors
Q = H(n-1) . . . H(2) H(1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(i+1:n) = 0 and v(i) = 1; v(1:i-1) is stored on exit in
A(1:i-1,i+1), and tau in TAU(i).
If UPLO = 'L', the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(n-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a complex scalar, and v is a complex vector with
v(1:i) = 0 and v(i+1) = 1; v(i+2:n) is stored on exit in A(i+2:n,i),
and tau in TAU(i).
The contents of A on exit are illustrated by the following examples
with n = 5:
if UPLO = 'U': if UPLO = 'L':
( d e v2 v3 v4 ) ( d )
( d e v3 v4 ) ( e d )
( d e v4 ) ( v1 e d )
( d e ) ( v1 v2 e d )
( d ) ( v1 v2 v3 e d )
where d and e denote diagonal and off-diagonal elements of T, and vi
denotes an element of the vector defining H(i).
===================================================================== */
#define a_ref(a_1,a_2) ( a + ((a_2)-1)*( lda) + (a_1)-1)
#define da_ref(a_1,a_2) (da + ((a_2)-1)*(ldda) + (a_1)-1)
#define tau_ref(a_1) (tau + (a_1)-1)
#define t_ref(a_1) (dT + ((a_1)-1)*(lddt))
#define Atest(a_1,a_2) ( Atest + ((a_2)-1)*( lda) + (a_1)-1)
#define dttest(a_0, a_1, a_2) (dTmgpu[a_0] + ((a_2)-1)*(lddt))
#define datest(a_0, a_1, a_2) (dAmgpu[a_0] + ((a_2)-1)*(ldda) + (a_1)-1)
char uplo_[2] = {uplo, 0};
magmaFloatComplex c_neg_one = MAGMA_C_NEG_ONE;
magmaFloatComplex c_neg_half = MAGMA_C_NEG_HALF;
magmaFloatComplex c_one = MAGMA_C_ONE ;
magmaFloatComplex c_zero = MAGMA_C_ZERO;
float d_one = MAGMA_D_ONE;
magma_int_t pm, pn, indi, indj, pk;
magma_int_t pm_old=0, pn_old=0, indi_old=0, indj_old=0, flipV=-1;
magma_int_t iblock, idev, di;
int i;
int lwkopt;
int lquery;
assert (nstream>=3);
assert (nstream>=(ngpu+1));
*info = 0;
int upper = lapackf77_lsame(uplo_, "U");
lquery = lwork == -1;
if (! upper && ! lapackf77_lsame(uplo_, "L")) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
} else if (lwork < 1 && ! lquery) {
*info = -9;
}
/* Determine the block size. */
lwkopt = n * nb;
if (*info == 0) {
MAGMA_C_SET2REAL( work[0], lwkopt );
}
if (*info != 0)
return *info;
else if (lquery)
return *info;
/* Quick return if possible */
if (n == 0) {
work[0] = c_one;
return *info;
}
magma_int_t mklth = min(threads,16);
magma_setlapack_numthreads(mklth);
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2];
magma_int_t nbcmplx=0;
magma_buildconnection_mgpu(gnode, &nbcmplx, ngpu);
#ifdef ENABLE_DEBUG
printf(" Initializing communication pattern.... GPU-ncmplx %d\n\n" , nbcmplx);
#endif
magma_device_t cdev;
magma_getdevice( &cdev );
magma_queue_t cstream;
magmablasGetKernelStream(&cstream);
magmaFloatComplex *dspace[MagmaMaxGPUs];
magmaFloatComplex *dwork[MagmaMaxGPUs], *dworkbis[MagmaMaxGPUs];
magmaFloatComplex *dvall[MagmaMaxGPUs], *dv[MagmaMaxGPUs], *dw[MagmaMaxGPUs];
magmaFloatComplex *workngpu[MagmaMaxGPUs+1];
magma_event_t redevents[MagmaMaxGPUs][MagmaMaxGPUs*MagmaMaxGPUs+10];
magma_int_t nbevents = MagmaMaxGPUs*MagmaMaxGPUs;
magma_int_t lddv = ldda;
magma_int_t lddw = lddv;
magma_int_t dwrk2siz = ldda*nb*(ngpu+1);
magma_int_t worksiz = n*nb;
magma_int_t devworksiz = 2*nb*lddv + nb*lddw + nb*ldda + dwrk2siz; // 2*dv(dv0+dv1) + dw + dwork +dworkbis
// local allocation and stream creation
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_cmalloc( &dspace[dev], devworksiz );
magma_cmalloc_pinned ( &workngpu[dev], worksiz);
dvall[dev] = dspace[dev];
dw[dev] = dvall[dev] + 2*nb*lddv;
dwork[dev] = dw[dev] + nb*lddw;
dworkbis[dev] = dwork[dev] + nb*ldda;
magmablasSetKernelStream( streams[ dev ][ 0 ] );
for( magma_int_t i = 0; i < nbevents; ++i ) {
cudaEventCreateWithFlags(&redevents[dev][i],cudaEventDisableTiming);
}
}
magma_cmalloc_pinned ( &workngpu[ngpu], worksiz);
magmaFloatComplex *worktest = NULL; //(magmaFloatComplex *) malloc(n*nb*sizeof(magmaFloatComplex)); // not used
// ======================
magmaFloatComplex *hT = work + lwork - nb*nb;
lwork -= nb*nb;
memset( hT, 0, nb*nb*sizeof(magmaFloatComplex));
if (upper) {
printf("CHETRD_HE2HB is not yet implemented for upper matrix storage. Exit.\n");
exit(1);
}else {
/* Reduce the lower triangle of A */
for (i = 1; i <= n-nb; i += nb)
{
indi = i+nb;
indj = i;
pm = n - i - nb + 1;
//pn = min(i+nb-1, n-nb) -i + 1;
pn = nb;
/* Get the current panel (no need for the 1st iteration) */
if (i > 1 ){
// cpanel_to_q copy the upper oof diagonal part of
// the matrix to work to be restored later. acctually
// the zero's and one's putted are not used this is only
// because we don't have a function that copy only the
// upper part of A to be restored after copying the
// lookahead panel that has been computted from GPU to CPU.
cpanel_to_q(MagmaUpper, pn-1, a_ref(i, i+1), lda, work);
// find the device who own the panel then send it to the CPU.
// below a -1 was added and then a -1 was done on di because of the fortran indexing
iblock = ((i-1) / distblk) / ngpu; // local block id
di = iblock*distblk + (i-1)%distblk; // local index in parent matrix
idev = ((i-1) / distblk) % ngpu; // device with this block
//printf("Receiving panel ofsize %d %d from idev %d A(%d,%d) \n",(pm+pn), pn,idev,i-1,di);
magma_setdevice( idev );
//magma_device_sync();
magma_cgetmatrix_async( (pm+pn), pn,
datest(idev, i, di+1), ldda,
a_ref ( i, i), lda, streams[ idev ][ nstream-1 ] );
/*
magma_device_sync();
cudaMemcpy2DAsync(a_ref(i,i), lda*sizeof(magmaFloatComplex),
datest(idev,i,di+1), ldda*sizeof(magmaFloatComplex),
(pm+pn)*sizeof(magmaFloatComplex), pn,
cudaMemcpyDeviceToHost, streams[ idev ][ nstream-1 ]);
*/
//magma_setdevice( 0 );
//printf("updating cher2k on A(%d,%d) of size %d %d \n",indi_old+pn_old-1,indi_old+pn_old-1,pm_old-pn_old,pn_old);
// compute CHER2K_MGPU
magmablas_cher2k_mgpu2(
MagmaLower, MagmaNoTrans, pm_old-pn_old, pn_old,
c_neg_one, dv, pm_old, pn_old,
dw, pm_old, pn_old,
d_one, dAmgpu, ldda, indi_old+pn_old-1,
ngpu, distblk, streams, 2 );
//magma_setdevice( 0 );
magma_setdevice( idev );
magma_queue_sync( streams[idev][ nstream-1 ] );
//magma_setdevice( 0 );
cq_to_panel(MagmaUpper, pn-1, a_ref(i, i+1), lda, work);
}
/* ==========================================================
QR factorization on a panel starting nb off of the diagonal.
Prepare the V and T matrices.
========================================================== */
lapackf77_cgeqrf(&pm, &pn, a_ref(indi, indj), &lda,
tau_ref(i), work, &lwork, info);
/* Form the matrix T */
pk=min(pm,pn);
lapackf77_clarft( MagmaForwardStr, MagmaColumnwiseStr,
&pm, &pk, a_ref(indi, indj), &lda,
tau_ref(i), hT, &nb);
/* Prepare V - put 0s in the upper triangular part of the panel
(and 1s on the diagonal), temporaly storing the original in work */
cpanel_to_q(MagmaUpper, pk, a_ref(indi, indj), lda, work);
/* Send V and T from the CPU to the GPU */
// To be able to overlap the GET with the CHER2K
// it should be done on last stream.
// TO Avoid a BUG that is overwriting the old_V
// used atthis moment by cher2k with the new_V
// send it now, we decide to have a flipflop
// vector of Vs. if step%2=0 use V[0] else use V[nb*n]
flipV = ((i-1)/nb)%2;
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
dv[dev] = dvall[dev] + flipV*nb*lddv;
}
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
// send V
magma_csetmatrix_async( pm, pk,
a_ref(indi, indj), lda,
dv[dev], pm, streams[dev][nstream-1] );
// Send the triangular factor T to the GPU
magma_csetmatrix_async( pk, pk,
hT, nb,
dttest(dev, 1, i), lddt, streams[dev][nstream-1] );
}
/* ==========================================================
Compute W:
1. X = A (V T)
2. W = X - 0.5* V * (T' * (V' * X))
========================================================== */
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
// dwork = V T
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ nstream-1 ] );
magma_queue_sync( streams[dev][nstream-1] );
magma_cgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_one, dv[dev], pm,
dttest(dev, 1, i), lddt,
c_zero, dwork[dev], pm);
}
// ===============================================
// SYNC TO BE SURE THAT BOTH V AND T WERE
// RECEIVED AND VT IS COMPUTED and SYR2K is done
// ===============================================
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
for( magma_int_t s = 0; s < nstream; ++s )
magma_queue_sync( streams[dev][s] );
}
// compute CHEMM_MGPU
// The broadcast of the result done inside this function
// should be done in stream [0] because i am assuming this
// for the GEMMs below otherwise I have to SYNC over the
// Broadcasting stream.
if(ngpu==1){
magmablasSetKernelStream( streams[ 0 ][ 0 ] );
magma_chemm(MagmaLeft, uplo, pm, pk,
c_one, dAmgpu[0]+(indi-1)*ldda+(indi-1), ldda,
dwork[0], pm,
c_zero, dw[0], pm);
}else{
magmablas_chemm_mgpu_com(
MagmaLeft, uplo, pm, pk,
c_one, dAmgpu, ldda, indi-1,
dwork, pm,
c_zero, dw, pm, dworkbis, dwrk2siz, worktest, pm, workngpu, worksiz,
ngpu, distblk, streams, nstream-1, redevents, nbevents, gnode, nbcmplx);
}
/* dwork = V*T already ==> dwork' = T'*V'
* compute T'*V'*X ==> dwork'*W ==>
* dwork + pm*nb = ((T' * V') * X) = dwork' * X = dwork' * W */
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
// Here we have to wait until the broadcast of CHEMM has been done.
// Note that the broadcast should be done on stream[0] so in a way
// we can continue here on the same stream and avoid a sync
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
// magma_queue_sync( streams[dev][0] );
magma_cgemm(MagmaConjTrans, MagmaNoTrans, pk, pk, pm,
c_one, dwork[dev], pm,
dw[dev], pm,
c_zero, dworkbis[dev], nb);
/* W = X - 0.5 * V * T'*V'*X
* = X - 0.5 * V * (dwork + pm*nb) = W - 0.5 * V * (dwork + pm*nb) */
magma_cgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_neg_half, dv[dev], pm,
dworkbis[dev], nb,
c_one, dw[dev], pm);
}
/* restore the panel it is put here to overlap with the previous GEMM*/
cq_to_panel(MagmaUpper, pk, a_ref(indi, indj), lda, work);
// ===============================================
// SYNC TO BE SURE THAT BOTH V AND W ARE DONE
// ===============================================
// Synchronise to be sure that W has been computed
// because next CHER2K use streaming and may happen
// that lunch a gemm on stream 2 while stream 0
// which compute those 2 GEMM above has not been
// computed and also used for the same reason in
// the panel update below and also for the last HER2K
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_queue_sync( streams[dev][0] );
}
/* ==========================================================
Update the unreduced submatrix A(i+ib:n,i+ib:n), using
an update of the form: A := A - V*W' - W*V'
========================================================== */
if (i + nb <= n-nb){
/* There would be next iteration;
do lookahead - update the next panel */
// below a -1 was added and then a -1 was done on di because of the fortran indexing
iblock = ((indi-1) / distblk) / ngpu; // local block id
di = iblock*distblk + (indi-1)%distblk; // local index in parent matrix
idev = ((indi-1) / distblk) % ngpu; // device with this block
magma_setdevice( idev );
magmablasSetKernelStream( streams[ idev ][ nstream-1 ] );
//magma_queue_sync( streams[idev][0] ); removed because the sync has been done in the loop above
magma_cgemm(MagmaNoTrans, MagmaConjTrans, pm, pn, pn, c_neg_one,
dv[idev], pm,
dw[idev] , pm, c_one,
datest(idev, indi, di+1), ldda);
magma_cgemm(MagmaNoTrans, MagmaConjTrans, pm, pn, pn, c_neg_one,
dw[idev] , pm,
dv[idev], pm, c_one,
datest(idev, indi, di+1), ldda);
//printf("updating next panel distblk %d idev %d on A(%d,%d) of size %d %d %d \n",distblk,idev,indi-1,di,pm,pn,pn);
}
else {
/* no look-ahead as this is last iteration */
// below a -1 was added and then a -1 was done on di because of the fortran indexing
iblock = ((indi-1) / distblk) / ngpu; // local block id
di = iblock*distblk + (indi-1)%distblk; // local index in parent matrix
idev = ((indi-1) / distblk) % ngpu; // device with this block
magma_setdevice( idev );
magmablasSetKernelStream( streams[ idev ][ 0 ] );
//printf("LAST CHER2K idev %d on A(%d,%d) of size %d \n",idev, indi-1,di,pk);
magma_cher2k(MagmaLower, MagmaNoTrans, pk, pk, c_neg_one,
dv[idev], pm,
dw[idev] , pm, d_one,
datest(idev, indi, di+1), ldda);
/* Send the last block to the CPU */
cpanel_to_q(MagmaUpper, pk-1, a_ref(n-pk+1, n-pk+2), lda, work);
magma_cgetmatrix( pk, pk,
datest(idev, indi, di+1), ldda,
a_ref(n-pk+1, n-pk+1), lda );
cq_to_panel(MagmaUpper, pk-1, a_ref(n-pk+1, n-pk+2), lda, work);
}
indi_old = indi;
indj_old = indj;
pm_old = pm;
pn_old = pn;
} // end loop for(i)
}// end of LOWER
//magma_setdevice( 0 );
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magma_free( dspace[dev]);
magma_free_pinned(workngpu[dev]);
for( magma_int_t e = 0; e < nbevents; ++e ) {
cudaEventDestroy(redevents[dev][e]);
}
}
magma_free_pinned(workngpu[ngpu]);
free(worktest);
magma_setdevice( cdev );
magmablasSetKernelStream( cstream );
MAGMA_C_SET2REAL( work[0], lwkopt );
magma_setlapack_numthreads(1);
return *info;
} /* chetrd_he2hb_ */
Event::Event( unsigned int flags )
: m_flags( flags )
{
DP_ASSERT( ( m_flags & ~( cudaEventDefault | cudaEventBlockingSync | cudaEventDisableTiming | cudaEventInterprocess ) ) == 0 )
CUDA_VERIFY( cudaEventCreateWithFlags( &m_event, m_flags ) );
}
explicit future(cudaStream_t s, bool owns_stream)
: m_stream(s),m_owns_stream(owns_stream)
{
detail::future_detail::throw_on_error(cudaEventCreateWithFlags(&m_event, create_flags), "cudaEventCreateWithFlags in future ctor");
detail::future_detail::throw_on_error(cudaEventRecord(m_event, m_stream), "cudaEventRecord in future ctor");
} // end future()
/**
Purpose
-------
DSYTRD_HE2HB reduces a real symmetric matrix A to real symmetric
band-diagonal form T by an orthogonal similarity transformation:
Q**H * A * Q = T.
This version stores the triangular matrices T used in the accumulated
Householder transformations (I - V T V').
Arguments
---------
@param[in]
uplo magma_uplo_t
- = MagmaUpper: Upper triangle of A is stored;
- = MagmaLower: Lower triangle of A is stored.
@param[in]
n INTEGER
The order of the matrix A. 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.
On exit, if UPLO = MagmaUpper, the Upper band-diagonal of A is
overwritten by the corresponding elements of the
band-diagonal matrix T, and the elements above the band
diagonal, with the array TAU, represent the orthogonal
matrix Q as a product of elementary reflectors; if UPLO
= MagmaLower, the the Lower band-diagonal of A is overwritten by
the corresponding elements of the band-diagonal
matrix T, and the elements below the band-diagonal, with
the array TAU, represent the orthogonal matrix Q as a product
of elementary reflectors. See Further Details.
@param[in]
lda INTEGER
The leading dimension of the array A. LDA >= max(1,N).
@param[out]
tau DOUBLE_PRECISION array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details).
@param[out]
work (workspace) DOUBLE_PRECISION array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK[0] returns the optimal LWORK.
@param[in]
lwork INTEGER
The dimension of the array WORK. LWORK >= 1.
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
\n
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
@param[out]
dT DOUBLE_PRECISION array on the GPU, dimension N*NB,
where NB is the optimal blocksize.
On exit dT holds the upper triangular matrices T from the
accumulated Householder transformations (I - V T V') used
in the factorization. The nb x nb matrices T are ordered
consecutively in memory one after another.
@param[out]
info INTEGER
- = 0: successful exit
- < 0: if INFO = -i, the i-th argument had an illegal value
Further Details
---------------
If UPLO = MagmaUpper, the matrix Q is represented as a product of elementary
reflectors
Q = H(n-1) . . . H(2) H(1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(i+1:n) = 0 and v(i) = 1; v(1:i-1) is stored on exit in
A(1:i-1,i+1), and tau in TAU(i).
If UPLO = MagmaLower, the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(n-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i) = 0 and v(i+1) = 1; v(i+2:n) is stored on exit in A(i+2:n,i),
and tau in TAU(i).
The contents of A on exit are illustrated by the following examples
with n = 5:
if UPLO = MagmaUpper: if UPLO = MagmaLower:
( d e v2 v3 v4 ) ( d )
( d e v3 v4 ) ( e d )
( d e v4 ) ( v1 e d )
( d e ) ( v1 v2 e d )
( d ) ( v1 v2 v3 e d )
where d and e denote diagonal and off-diagonal elements of T, and vi
denotes an element of the vector defining H(i).
@ingroup magma_dsyev_2stage
********************************************************************/
extern "C" magma_int_t
magma_dsytrd_sy2sb( magma_uplo_t uplo, magma_int_t n, magma_int_t nb,
double *A, magma_int_t lda,
double *tau,
double *work, magma_int_t lwork,
double *dT,
magma_int_t *info)
{
#define A(a_1,a_2) ( A + ((a_2)-1)*( lda) + (a_1)-1)
#define dA(a_1,a_2) (dA + ((a_2)-1)*(ldda) + (a_1)-1)
#define tau_ref(a_1) (tau + (a_1)-1)
#define dT(a_1) (dT + ((a_1)-1)*(lddt))
int ldda = ((n+31)/32)*32;
int lddt = nb;
double c_neg_one = MAGMA_D_NEG_ONE;
double c_neg_half = MAGMA_D_NEG_HALF;
double c_one = MAGMA_D_ONE;
double c_zero = MAGMA_D_ZERO;
double d_one = MAGMA_D_ONE;
magma_int_t pm, pn, indi, indj, pk;
magma_int_t pm_old=0, pn_old=0, indi_old=0, indj_old=0;
int i;
int lwkopt;
int lquery;
*info = 0;
int upper = (uplo == MagmaUpper);
lquery = (lwork == -1);
if (! upper && uplo != MagmaLower) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
} else if (lwork < 1 && ! lquery) {
*info = -9;
}
/* Determine the block size. */
lwkopt = n * nb;
if (*info == 0) {
work[0] = MAGMA_D_MAKE( lwkopt, 0 );
}
if (*info != 0)
return *info;
else if (lquery)
return *info;
/* Quick return if possible */
if (n == 0) {
work[0] = c_one;
return *info;
}
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
double *dA;
if (MAGMA_SUCCESS != magma_dmalloc( &dA, (n + 2*nb)*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
// limit to 16 threads
magma_int_t orig_threads = magma_get_lapack_numthreads();
magma_set_lapack_numthreads( min(orig_threads,16) );
/* Use the first panel of dA as work space */
double *dwork = dA + n*ldda;
double *dW = dwork + nb*ldda;
#ifdef TRACING
char buf[80];
#endif
magma_queue_t stream[3];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
stream[2] = 0; // default stream
trace_init( 1, 1, 3, stream );
double *hT = work + lwork - nb*nb;
lwork -= nb*nb;
memset( hT, 0, nb*nb*sizeof(double));
magmablasSetKernelStream( stream[0] );
magma_event_t Pupdate_event;
cudaEventCreateWithFlags(&Pupdate_event,cudaEventDisableTiming);
//magma_event_create(&Pupdate_event);
if (upper) {
printf("DSYTRD_HE2HB is not yet implemented for upper matrix storage. Exit.\n");
exit(1);
} else {
/* Copy the matrix to the GPU */
if (1 <= n-nb) {
trace_gpu_start( 0, 0, "set", "set A" );
magma_dsetmatrix_async( (n-nb), (n-nb),
A(nb+1, nb+1), lda,
dA(nb+1, nb+1), ldda, stream[0] );
trace_gpu_end( 0, 0 );
}
/* Reduce the lower triangle of A */
for (i = 1; i <= n-nb; i += nb) {
indi = i+nb;
indj = i;
pm = n - i - nb + 1;
//pn = min(i+nb-1, n-nb) -i + 1;
pn = nb;
/* Get the current panel (no need for the 1st iteration) */
if (i > 1 ) {
// dpanel_to_q copy the upper oof diagonal part of
// the matrix to work to be restored later. acctually
// the zero's and one's putted are not used this is only
// because we don't have a function that copy only the
// upper part of A to be restored after copying the
// lookahead panel that has been computted from GPU to CPU.
dpanel_to_q(MagmaUpper, pn-1, A(i, i+1), lda, work);
trace_gpu_start( 0, 1, "get", "get panel" );
//magma_queue_sync( stream[0] );
magma_queue_wait_event(stream[1], Pupdate_event); //, 0);
magma_dgetmatrix_async( (pm+pn), pn,
dA( i, i), ldda,
A ( i, i), lda, stream[1] );
trace_gpu_end( 0, 1 );
trace_gpu_start( 0, 2, "her2k", "her2k" );
magma_dsyr2k(MagmaLower, MagmaNoTrans, pm_old-pn_old, pn_old, c_neg_one,
dA(indi_old+pn_old, indj_old), ldda,
dW + pn_old, pm_old, d_one,
dA(indi_old+pn_old, indi_old+pn_old), ldda);
trace_gpu_end( 0, 2 );
trace_cpu_start( 0, "sync", "sync on 1" );
magma_queue_sync( stream[1] );
trace_cpu_end( 0 );
dq_to_panel(MagmaUpper, pn-1, A(i, i+1), lda, work);
}
/* ==========================================================
QR factorization on a panel starting nb off of the diagonal.
Prepare the V and T matrices.
========================================================== */
#ifdef TRACING
snprintf( buf, sizeof(buf), "panel %d", i );
#endif
trace_cpu_start( 0, "geqrf", buf );
lapackf77_dgeqrf(&pm, &pn, A(indi, indj), &lda,
tau_ref(i), work, &lwork, info);
/* Form the matrix T */
pk=min(pm,pn);
lapackf77_dlarft( MagmaForwardStr, MagmaColumnwiseStr,
&pm, &pk, A(indi, indj), &lda,
tau_ref(i), hT, &nb);
/* Prepare V - put 0s in the upper triangular part of the panel
(and 1s on the diagonal), temporaly storing the original in work */
dpanel_to_q(MagmaUpper, pk, A(indi, indj), lda, work);
trace_cpu_end( 0 );
/* Send V from the CPU to the GPU */
trace_gpu_start( 0, 0, "set", "set V and T" );
magma_dsetmatrix_async( pm, pk,
A(indi, indj), lda,
dA(indi, indj), ldda, stream[0] );
/* Send the triangular factor T to the GPU */
magma_dsetmatrix_async( pk, pk,
hT, nb,
dT(i), lddt, stream[0] );
trace_gpu_end( 0, 0 );
/* ==========================================================
Compute W:
1. X = A (V T)
2. W = X - 0.5* V * (T' * (V' * X))
========================================================== */
/* dwork = V T */
trace_cpu_start( 0, "sync", "sync on 0" );
// this sync is done here to be sure that the copy has been finished
// because below we made a restore dq_to_panel and this restore need
// to ensure that the copy has been finished. we did it here to allow
// overlapp of restore with next gemm and symm.
magma_queue_sync( stream[0] );
trace_cpu_end( 0 );
trace_gpu_start( 0, 2, "gemm", "work = V*T" );
magma_dgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_one, dA(indi, indj), ldda,
dT(i), lddt,
c_zero, dwork, pm);
trace_gpu_end( 0, 2 );
/* dW = X = A*V*T. dW = A*dwork */
trace_gpu_start( 0, 2, "hemm", "X = A*work" );
magma_dsymm(MagmaLeft, uplo, pm, pk,
c_one, dA(indi, indi), ldda,
dwork, pm,
c_zero, dW, pm);
trace_gpu_end( 0, 2 );
/* restore the panel */
dq_to_panel(MagmaUpper, pk, A(indi, indj), lda, work);
/* dwork = V*T already ==> dwork' = T'*V'
* compute T'*V'*X ==> dwork'*W ==>
* dwork + pm*nb = ((T' * V') * X) = dwork' * X = dwork' * W */
trace_gpu_start( 0, 2, "gemm", "work = T'*V'*X" );
magma_dgemm(MagmaConjTrans, MagmaNoTrans, pk, pk, pm,
c_one, dwork, pm,
dW, pm,
c_zero, dwork + pm*nb, nb);
trace_gpu_end( 0, 2 );
/* W = X - 0.5 * V * T'*V'*X
* = X - 0.5 * V * (dwork + pm*nb) = W - 0.5 * V * (dwork + pm*nb) */
trace_gpu_start( 0, 2, "gemm", "W = X - 0.5*V*(T'*V'*X)" );
magma_dgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_neg_half, dA(indi, indj), ldda,
dwork + pm*nb, nb,
c_one, dW, pm);
trace_gpu_end( 0, 2 );
/* ==========================================================
Update the unreduced submatrix A(i+ib:n,i+ib:n), using
an update of the form: A := A - V*W' - W*V'
========================================================== */
if (i + nb <= n-nb) {
/* There would be next iteration;
do lookahead - update the next panel */
trace_gpu_start( 0, 2, "gemm", "gemm 4 next panel left" );
magma_dgemm(MagmaNoTrans, MagmaConjTrans, pm, pn, pn, c_neg_one,
dA(indi, indj), ldda,
dW, pm, c_one,
dA(indi, indi), ldda);
trace_gpu_end( 0, 2 );
trace_gpu_start( 0, 2, "gemm", "gemm 5 next panel right" );
magma_dgemm(MagmaNoTrans, MagmaConjTrans, pm, pn, pn, c_neg_one,
dW, pm,
dA(indi, indj), ldda, c_one,
dA(indi, indi), ldda);
trace_gpu_end( 0, 2 );
magma_event_record(Pupdate_event, stream[0]);
}
else {
/* no look-ahead as this is last iteration */
trace_gpu_start( 0, 2, "her2k", "her2k last iteration" );
magma_dsyr2k(MagmaLower, MagmaNoTrans, pk, pk, c_neg_one,
dA(indi, indj), ldda,
dW, pm, d_one,
dA(indi, indi), ldda);
trace_gpu_end( 0, 2 );
}
indi_old = indi;
indj_old = indj;
pm_old = pm;
pn_old = pn;
} // end loop for (i)
/* Send the last block to the CPU */
pk = min(pm,pn);
if (1 <= n-nb) {
dpanel_to_q(MagmaUpper, pk-1, A(n-pk+1, n-pk+2), lda, work);
trace_gpu_start( 0, 2, "get", "get last block" );
magma_dgetmatrix( pk, pk,
dA(n-pk+1, n-pk+1), ldda,
A(n-pk+1, n-pk+1), lda );
trace_gpu_end( 0, 2 );
dq_to_panel(MagmaUpper, pk-1, A(n-pk+1, n-pk+2), lda, work);
}
}// end of LOWER
trace_finalize( "dsytrd_sy2sb.svg", "trace.css" );
magma_event_destroy( Pupdate_event );
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free( dA );
work[0] = MAGMA_D_MAKE( lwkopt, 0 );
magmablasSetKernelStream( orig_stream );
magma_set_lapack_numthreads( orig_threads );
return *info;
} /* magma_dsytrd_sy2sb */
gaspi_return_t
gaspi_init_GPUs()
{
int i, j, k;
int deviceCount;
int device_id = 0;
int gaspi_devices = 0;
int ib_numa_node;
int direct_devices[32];
struct cudaDeviceProp deviceProp;
cudaError_t error_id = cudaGetDeviceCount(&deviceCount);
if ( error_id != cudaSuccess )
{
gaspi_print_error("Failed cudaGetDeviceCount." );
return GASPI_ERROR;
}
if( deviceCount <= 0 )
{
gaspi_print_error("No CUDA capable devices found.");
return GASPI_ERROR;
}
ib_numa_node = _gaspi_find_GPU_ib_numa_node();
for(device_id = 0; device_id < deviceCount; device_id++)
{
cudaGetDeviceProperties(&deviceProp, device_id);
if( deviceProp.major >= 3 ) /* TODO: magic number */
{
cudaSetDevice(device_id);
if( ib_numa_node == _gaspi_find_GPU_numa_node(device_id) )
{
direct_devices[gaspi_devices] = device_id;
gaspi_devices++;
}
}
}
if( 0 == gaspi_devices )
{
gaspi_print_error("No GPU Direct RDMA capable devices on the correct NUMA-socket were found.");
return GASPI_ERROR;
}
glb_gaspi_ctx.gpu_count = gaspi_devices;
gpus = (gaspi_gpu *) malloc(sizeof(gaspi_gpu)*glb_gaspi_ctx.gpu_count);
if( !gpus )
{
gaspi_print_error("Failed to allocate mameory.");
return GASPI_ERR_MEMALLOC;
}
for(k = 0 ; k < gaspi_devices; k++)
{
cudaSetDevice(direct_devices[k]);
for( i = 0; i < GASPI_MAX_QP; i++)
{
cudaStreamCreate(&gpus[k].streams[i]);
for(j = 0; j < GASPI_CUDA_EVENTS; j++)
{
cudaEventCreateWithFlags(&gpus[k].events[i][j].event, cudaEventDisableTiming);
}
cudaStreamCreateWithFlags(&gpus[k].streams[i], cudaStreamNonBlocking);
}
gpus[k].device_id = direct_devices[k];
}
glb_gaspi_ctx.use_gpus = 1;
return GASPI_SUCCESS;
}
magma_int_t
magma_d_initP2P ( magma_int_t *bw_bmark, magma_int_t *num_gpus ){
// Number of GPUs
printf("Checking for multiple GPUs...\n");
int gpu_n;
(cudaGetDeviceCount(&gpu_n));
printf("CUDA-capable device count: %i\n", gpu_n);
if (gpu_n < 2)
{
printf("Two or more Tesla(s) with (SM 2.0)"
" class GPUs are required for P2P.\n");
}
// Query device properties
cudaDeviceProp prop[64];
int gpuid_tesla[64]; // find the first two GPU's that can support P2P
int gpu_count = 0; // GPUs that meet the criteria
for (int i=0; i < gpu_n; i++) {
(cudaGetDeviceProperties(&prop[i], i));
// Only Tesla boards based on Fermi can support P2P
{
// This is an array of P2P capable GPUs
gpuid_tesla[gpu_count++] = i;
}
}
*num_gpus=gpu_n;
for(int i=0; i<gpu_n; i++)
{
for(int j=i+1; j<gpu_n; j++)
{
// Check possibility for peer access
printf("\nChecking GPU(s) for support of peer to peer memory access...\n");
int can_access_peer_0_1, can_access_peer_1_0;
// In this case we just pick the first two that we can support
(cudaDeviceCanAccessPeer(&can_access_peer_0_1, gpuid_tesla[i],
gpuid_tesla[j]));
(cudaDeviceCanAccessPeer(&can_access_peer_1_0, gpuid_tesla[j],
gpuid_tesla[i]));
// Output results from P2P capabilities
printf("> Peer access from %s (GPU%d) -> %s (GPU%d) : %s\n",
prop[gpuid_tesla[i]].name, gpuid_tesla[i],
prop[gpuid_tesla[j]].name, gpuid_tesla[j] ,
can_access_peer_0_1 ? "Yes" : "No");
printf("> Peer access from %s (GPU%d) -> %s (GPU%d) : %s\n",
prop[gpuid_tesla[j]].name, gpuid_tesla[j],
prop[gpuid_tesla[i]].name, gpuid_tesla[i],
can_access_peer_1_0 ? "Yes" : "No");
if (can_access_peer_0_1 == 0 || can_access_peer_1_0 == 0)
{
printf("Two or more Tesla(s) with class"
" GPUs are required for P2P to run.\n");
printf("Support for UVA requires a Tesla with SM 2.0 capabilities.\n");
printf("Peer to Peer access is not available between"
" GPU%d <-> GPU%d, waiving test.\n", gpuid_tesla[i], gpuid_tesla[j]);
printf("PASSED\n");
//exit(EXIT_SUCCESS);
}
}
}
// Enable peer access
for(int i=0; i<gpu_n; i++)
{
for(int j=i+1; j<gpu_n; j++)
{
printf("Enabling peer access between GPU%d and GPU%d...\n",
gpuid_tesla[i], gpuid_tesla[j]);
(cudaSetDevice(gpuid_tesla[i]));
(cudaDeviceEnablePeerAccess(gpuid_tesla[j], 0));
(cudaSetDevice(gpuid_tesla[j]));
(cudaDeviceEnablePeerAccess(gpuid_tesla[i], 0));
magma_dcheckerr("P2P");
}
}
magma_dcheckerr("P2P successful");
// Enable peer access
for(int i=0; i<gpu_n; i++)
{
for(int j=i+1; j<gpu_n; j++)
{
// Check that we got UVA on both devices
printf("Checking GPU%d and GPU%d for UVA capabilities...\n",
gpuid_tesla[i], gpuid_tesla[j]);
//const bool has_uva = (prop[gpuid_tesla[i]].unifiedAddressing &&
// prop[gpuid_tesla[j]].unifiedAddressing);
printf("> %s (GPU%d) supports UVA: %s\n", prop[gpuid_tesla[i]].name,
gpuid_tesla[i], (prop[gpuid_tesla[i]].unifiedAddressing ? "Yes" : "No") );
printf("> %s (GPU%d) supports UVA: %s\n", prop[gpuid_tesla[j]].name,
gpuid_tesla[j], (prop[gpuid_tesla[j]].unifiedAddressing ? "Yes" : "No") );
}
}
if(*bw_bmark==1){
// P2P memcopy() benchmark
for(int i=0; i<gpu_n; i++)
{
for(int j=i+1; j<gpu_n; j++)
{
// Allocate buffers
const size_t buf_size = 1024 * 1024 * 16 * sizeof(float);
printf("Allocating buffers (%iMB on GPU%d, GPU%d and CPU Host)...\n",
int(buf_size / 1024 / 1024), gpuid_tesla[i], gpuid_tesla[j]);
(cudaSetDevice(gpuid_tesla[i]));
float* g0;
(cudaMalloc(&g0, buf_size));
(cudaSetDevice(gpuid_tesla[j]));
float* g1;
(cudaMalloc(&g1, buf_size));
float* h0;
(cudaMallocHost(&h0, buf_size)); // Automatically portable with UVA
// Create CUDA event handles
printf("Creating event handles...\n");
cudaEvent_t start_event, stop_event;
float time_memcpy;
int eventflags = cudaEventBlockingSync;
(cudaEventCreateWithFlags(&start_event, eventflags));
(cudaEventCreateWithFlags(&stop_event, eventflags));
(cudaEventRecord(start_event, 0));
for (int k=0; k<100; k++)
{
// With UVA we don't need to specify source and target devices, the
// runtime figures this out by itself from the pointers
// Ping-pong copy between GPUs
if (k % 2 == 0)
(cudaMemcpy(g1, g0, buf_size, cudaMemcpyDefault));
else
(cudaMemcpy(g0, g1, buf_size, cudaMemcpyDefault));
}
(cudaEventRecord(stop_event, 0));
(cudaEventSynchronize(stop_event));
(cudaEventElapsedTime(&time_memcpy, start_event, stop_event));
printf("cudaMemcpyPeer / cudaMemcpy between"
"GPU%d and GPU%d: %.2fGB/s\n", gpuid_tesla[i], gpuid_tesla[j],
(1.0f / (time_memcpy / 1000.0f))
* ((100.0f * buf_size)) / 1024.0f / 1024.0f / 1024.0f);
// Cleanup and shutdown
printf("Cleanup of P2P benchmark...\n");
(cudaEventDestroy(start_event));
(cudaEventDestroy(stop_event));
(cudaSetDevice(gpuid_tesla[i]));
(magma_free( g0) );
(cudaSetDevice(gpuid_tesla[j]));
(magma_free( g1) );
(magma_free_cpu( h0) );
}
}
// host-device memcopy() benchmark
for(int j=0; j<gpu_n; j++)
{
cudaSetDevice(gpuid_tesla[j]);
int *h_data_source;
int *h_data_sink;
int *h_data_in[STREAM_COUNT];
int *d_data_in[STREAM_COUNT];
int *h_data_out[STREAM_COUNT];
int *d_data_out[STREAM_COUNT];
cudaEvent_t cycleDone[STREAM_COUNT];
cudaStream_t stream[STREAM_COUNT];
cudaEvent_t start, stop;
// Allocate resources
int memsize;
memsize = 1000000 * sizeof(int);
h_data_source = (int*) malloc(memsize);
h_data_sink = (int*) malloc(memsize);
for( int i =0; i<STREAM_COUNT; ++i ) {
( cudaHostAlloc(&h_data_in[i], memsize,
cudaHostAllocDefault) );
( cudaMalloc(&d_data_in[i], memsize) );
( cudaHostAlloc(&h_data_out[i], memsize,
cudaHostAllocDefault) );
( cudaMalloc(&d_data_out[i], memsize) );
( cudaStreamCreate(&stream[i]) );
( cudaEventCreate(&cycleDone[i]) );
cudaEventRecord(cycleDone[i], stream[i]);
}
cudaEventCreate(&start); cudaEventCreate(&stop);
// Time host-device copies
cudaEventRecord(start,0);
( cudaMemcpyAsync(d_data_in[0], h_data_in[0], memsize,
cudaMemcpyHostToDevice,0) );
cudaEventRecord(stop,0);
cudaEventSynchronize(stop);
float memcpy_h2d_time;
cudaEventElapsedTime(&memcpy_h2d_time, start, stop);
cudaEventRecord(start,0);
( cudaMemcpyAsync(h_data_out[0], d_data_out[0], memsize,
cudaMemcpyDeviceToHost, 0) );
cudaEventRecord(stop,0);
cudaEventSynchronize(stop);
float memcpy_d2h_time;
cudaEventElapsedTime(&memcpy_d2h_time, start, stop);
cudaEventSynchronize(stop);
printf("Measured timings (throughput):\n");
printf(" Memcpy host to device GPU %d \t: %f ms (%f GB/s)\n", j,
memcpy_h2d_time, (memsize * 1e-6)/ memcpy_h2d_time );
printf(" Memcpy device GPU %d to host\t: %f ms (%f GB/s)\n", j,
memcpy_d2h_time, (memsize * 1e-6)/ memcpy_d2h_time);
// Free resources
free( h_data_source );
free( h_data_sink );
for( int i =0; i<STREAM_COUNT; ++i ) {
magma_free_cpu( h_data_in[i] );
magma_free( d_data_in[i] );
magma_free_cpu( h_data_out[i] );
magma_free( d_data_out[i] );
cudaStreamDestroy(stream[i]);
cudaEventDestroy(cycleDone[i]);
}
cudaEventDestroy(start);
cudaEventDestroy(stop);
}
}//end if-loop bandwidth_benchmark
magma_dcheckerr("P2P established");
return MAGMA_SUCCESS;
}
extern "C" magma_int_t
magma_ssytrd_sy2sb( char uplo, magma_int_t n, magma_int_t nb,
float *a, magma_int_t lda,
float *tau,
float *work, magma_int_t lwork,
float *dT,
magma_int_t threads, magma_int_t *info)
{
/* -- MAGMA (version 1.3.0) --
Univ. of Tennessee, Knoxville
Univ. of California, Berkeley
Univ. of Colorado, Denver
November 2012
Purpose
=======
SSYTRD_HE2HB reduces a real symmetric matrix A to real symmetric
band-diagonal form T by an orthogonal similarity transformation:
Q**T * A * Q = T.
This version stores the triangular matrices T used in the accumulated
Householder transformations (I - V T V').
Arguments
=========
UPLO (input) CHARACTER*1
= 'U': Upper triangle of A is stored;
= 'L': Lower triangle of A is stored.
N (input) INTEGER
The order of the matrix A. N >= 0.
A (input/output) REAL array, dimension (LDA,N)
On entry, the symmetric 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 UPLO = 'U', the Upper band-diagonal of A is
overwritten by the corresponding elements of the
band-diagonal matrix T, and the elements above the band
diagonal, with the array TAU, represent the orthogonal
matrix Q as a product of elementary reflectors; if UPLO
= 'L', the the Lower band-diagonal of A is overwritten by
the corresponding elements of the band-diagonal
matrix T, and the elements below the band-diagonal, with
the array TAU, represent the orthogonal matrix Q as a product
of elementary reflectors. See Further Details.
LDA (input) INTEGER
The leading dimension of the array A. LDA >= max(1,N).
TAU (output) REAL array, dimension (N-1)
The scalar factors of the elementary reflectors (see Further
Details).
WORK (workspace/output) REAL array, dimension (MAX(1,LWORK))
On exit, if INFO = 0, WORK(1) returns the optimal LWORK.
LWORK (input) INTEGER
The dimension of the array WORK. LWORK >= 1.
For optimum performance LWORK >= N*NB, where NB is the
optimal blocksize.
If LWORK = -1, then a workspace query is assumed; the routine
only calculates the optimal size of the WORK array, returns
this value as the first entry of the WORK array, and no error
message related to LWORK is issued by XERBLA.
dT (output) REAL array on the GPU, dimension N*NB,
where NB is the optimal blocksize.
On exit dT holds the upper triangular matrices T from the
accumulated Householder transformations (I - V T V') used
in the factorization. The nb x nb matrices T are ordered
consecutively in memory one after another.
INFO (output) INTEGER
= 0: successful exit
< 0: if INFO = -i, the i-th argument had an illegal value
Further Details
===============
If UPLO = 'U', the matrix Q is represented as a product of elementary
reflectors
Q = H(n-1) . . . H(2) H(1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(i+1:n) = 0 and v(i) = 1; v(1:i-1) is stored on exit in
A(1:i-1,i+1), and tau in TAU(i).
If UPLO = 'L', the matrix Q is represented as a product of elementary
reflectors
Q = H(1) H(2) . . . H(n-1).
Each H(i) has the form
H(i) = I - tau * v * v'
where tau is a real scalar, and v is a real vector with
v(1:i) = 0 and v(i+1) = 1; v(i+2:n) is stored on exit in A(i+2:n,i),
and tau in TAU(i).
The contents of A on exit are illustrated by the following examples
with n = 5:
if UPLO = 'U': if UPLO = 'L':
( d e v2 v3 v4 ) ( d )
( d e v3 v4 ) ( e d )
( d e v4 ) ( v1 e d )
( d e ) ( v1 v2 e d )
( d ) ( v1 v2 v3 e d )
where d and e denote diagonal and off-diagonal elements of T, and vi
denotes an element of the vector defining H(i).
===================================================================== */
#define a_ref(a_1,a_2) ( a + ((a_2)-1)*( lda) + (a_1)-1)
#define da_ref(a_1,a_2) (da + ((a_2)-1)*(ldda) + (a_1)-1)
#define tau_ref(a_1) (tau + (a_1)-1)
#define t_ref(a_1) (dT + ((a_1)-1)*(lddt))
char uplo_[2] = {uplo, 0};
int ldda = ((n+31)/32)*32;
int lddt = nb;
float c_neg_one = MAGMA_S_NEG_ONE;
float c_neg_half = MAGMA_S_NEG_HALF;
float c_one = MAGMA_S_ONE ;
float c_zero = MAGMA_S_ZERO;
float d_one = MAGMA_D_ONE;
magma_int_t pm, pn, indi, indj, pk;
magma_int_t pm_old=0, pn_old=0, indi_old=0, indj_old=0;
int i;
int lwkopt;
int lquery;
*info = 0;
int upper = lapackf77_lsame(uplo_, "U");
lquery = lwork == -1;
if (! upper && ! lapackf77_lsame(uplo_, "L")) {
*info = -1;
} else if (n < 0) {
*info = -2;
} else if (lda < max(1,n)) {
*info = -4;
} else if (lwork < 1 && ! lquery) {
*info = -9;
}
if (*info == 0) {
/* Determine the block size. */
lwkopt = n * nb;
MAGMA_S_SET2REAL( work[0], lwkopt );
}
if (*info != 0)
return *info;
else if (lquery)
return *info;
/* Quick return if possible */
if (n == 0) {
work[0] = c_one;
return *info;
}
float *da;
if (MAGMA_SUCCESS != magma_smalloc( &da, (n + 2*nb)*ldda )) {
*info = MAGMA_ERR_DEVICE_ALLOC;
return *info;
}
magma_int_t mklth = min(threads,12);
#if defined(USEMKL)
mkl_set_num_threads(mklth);
#endif
#if defined(USEACML)
omp_set_num_threads(mklth);
#endif
/* Use the first panel of da as work space */
float *dwork = da+n*ldda;
float *dW = dwork + nb*ldda;
#ifdef TRACING
char buf[80];
#endif
cudaStream_t stream[3];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
stream[2] = 0; // default stream
trace_init( 1, 1, 3, stream );
float *hT = work + lwork - nb*nb;
lwork -= nb*nb;
memset( hT, 0, nb*nb*sizeof(float));
magmablasSetKernelStream( stream[0] );
cudaEvent_t Pupdate_event;
cudaEventCreateWithFlags(&Pupdate_event,cudaEventDisableTiming);
//cudaEventCreate(&Pupdate_event);
if (upper) {
printf("SSYTRD_HE2HB is not yet implemented for upper matrix storage. Exit.\n");
exit(1);
}else {
/* Copy the matrix to the GPU */
if (1 <= n-nb){
trace_gpu_start( 0, 0, "set", "set A" );
magma_ssetmatrix_async( (n-nb), (n-nb),
a_ref(nb+1, nb+1), lda,
da_ref(nb+1, nb+1), ldda, stream[0] );
trace_gpu_end( 0, 0 );
}
/* Reduce the lower triangle of A */
for (i = 1; i <= n-nb; i += nb)
{
indi = i+nb;
indj = i;
pm = n - i - nb + 1;
//pn = min(i+nb-1, n-nb) -i + 1;
pn = nb;
/* Get the current panel (no need for the 1st iteration) */
if (i > 1 ){
// spanel_to_q copy the upper oof diagonal part of
// the matrix to work to be restored later. acctually
// the zero's and one's putted are not used this is only
// because we don't have a function that copy only the
// upper part of A to be restored after copying the
// lookahead panel that has been computted from GPU to CPU.
spanel_to_q(MagmaUpper, pn-1, a_ref(i, i+1), lda, work);
trace_gpu_start( 0, 1, "get", "get panel" );
//magma_queue_sync( stream[0] );
cudaStreamWaitEvent(stream[1], Pupdate_event, 0);
magma_sgetmatrix_async( (pm+pn), pn,
da_ref( i, i), ldda,
a_ref ( i, i), lda, stream[1] );
trace_gpu_end( 0, 1 );
trace_gpu_start( 0, 2, "syr2k", "syr2k" );
magma_ssyr2k(MagmaLower, MagmaNoTrans, pm_old-pn_old, pn_old, c_neg_one,
da_ref(indi_old+pn_old, indj_old), ldda,
dW + pn_old , pm_old, d_one,
da_ref(indi_old+pn_old, indi_old+pn_old), ldda);
trace_gpu_end( 0, 2 );
trace_cpu_start( 0, "sync", "sync on 1" );
magma_queue_sync( stream[1] );
trace_cpu_end( 0 );
sq_to_panel(MagmaUpper, pn-1, a_ref(i, i+1), lda, work);
}
/* ==========================================================
QR factorization on a panel starting nb off of the diagonal.
Prepare the V and T matrices.
========================================================== */
#ifdef TRACING
snprintf( buf, sizeof(buf), "panel %d", i );
#endif
trace_cpu_start( 0, "geqrf", buf );
lapackf77_sgeqrf(&pm, &pn, a_ref(indi, indj), &lda,
tau_ref(i), work, &lwork, info);
/* Form the matrix T */
pk=min(pm,pn);
lapackf77_slarft( MagmaForwardStr, MagmaColumnwiseStr,
&pm, &pk, a_ref(indi, indj), &lda,
tau_ref(i), hT, &nb);
/* Prepare V - put 0s in the upper triangular part of the panel
(and 1s on the diagonal), temporaly storing the original in work */
spanel_to_q(MagmaUpper, pk, a_ref(indi, indj), lda, work);
trace_cpu_end( 0 );
/* Send V from the CPU to the GPU */
trace_gpu_start( 0, 0, "set", "set V and T" );
magma_ssetmatrix_async( pm, pk,
a_ref(indi, indj), lda,
da_ref(indi, indj), ldda, stream[0] );
/* Send the triangular factor T to the GPU */
magma_ssetmatrix_async( pk, pk,
hT, nb,
t_ref(i), lddt, stream[0] );
trace_gpu_end( 0, 0 );
/* ==========================================================
Compute W:
1. X = A (V T)
2. W = X - 0.5* V * (T' * (V' * X))
========================================================== */
/* dwork = V T */
trace_cpu_start( 0, "sync", "sync on 0" );
// this sync is done here to be sure that the copy has been finished
// because below we made a restore sq_to_panel and this restore need
// to ensure that the copy has been finished. we did it here to allow
// overlapp of restore with next gemm and symm.
magma_queue_sync( stream[0] );
trace_cpu_end( 0 );
trace_gpu_start( 0, 2, "gemm", "work = V*T" );
magma_sgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_one, da_ref(indi, indj), ldda,
t_ref(i), lddt,
c_zero, dwork, pm);
trace_gpu_end( 0, 2 );
/* dW = X = A*V*T. dW = A*dwork */
trace_gpu_start( 0, 2, "symm", "X = A*work" );
magma_ssymm(MagmaLeft, uplo, pm, pk,
c_one, da_ref(indi, indi), ldda,
dwork, pm,
c_zero, dW, pm);
trace_gpu_end( 0, 2 );
/* restore the panel */
sq_to_panel(MagmaUpper, pk, a_ref(indi, indj), lda, work);
/* dwork = V*T already ==> dwork' = T'*V'
* compute T'*V'*X ==> dwork'*W ==>
* dwork + pm*nb = ((T' * V') * X) = dwork' * X = dwork' * W */
trace_gpu_start( 0, 2, "gemm", "work = T'*V'*X" );
magma_sgemm(MagmaTrans, MagmaNoTrans, pk, pk, pm,
c_one, dwork, pm,
dW, pm,
c_zero, dwork + pm*nb, nb);
trace_gpu_end( 0, 2 );
/* W = X - 0.5 * V * T'*V'*X
* = X - 0.5 * V * (dwork + pm*nb) = W - 0.5 * V * (dwork + pm*nb) */
trace_gpu_start( 0, 2, "gemm", "W = X - 0.5*V*(T'*V'*X)" );
magma_sgemm(MagmaNoTrans, MagmaNoTrans, pm, pk, pk,
c_neg_half, da_ref(indi, indj), ldda,
dwork + pm*nb, nb,
c_one, dW, pm);
trace_gpu_end( 0, 2 );
/* ==========================================================
Update the unreduced submatrix A(i+ib:n,i+ib:n), using
an update of the form: A := A - V*W' - W*V'
========================================================== */
if (i + nb <= n-nb){
/* There would be next iteration;
do lookahead - update the next panel */
trace_gpu_start( 0, 2, "gemm", "gemm 4 next panel left" );
magma_sgemm(MagmaNoTrans, MagmaTrans, pm, pn, pn, c_neg_one,
da_ref(indi, indj), ldda,
dW , pm, c_one,
da_ref(indi, indi), ldda);
trace_gpu_end( 0, 2 );
trace_gpu_start( 0, 2, "gemm", "gemm 5 next panel right" );
magma_sgemm(MagmaNoTrans, MagmaTrans, pm, pn, pn, c_neg_one,
dW , pm,
da_ref(indi, indj), ldda, c_one,
da_ref(indi, indi), ldda);
trace_gpu_end( 0, 2 );
cudaEventRecord(Pupdate_event, stream[0]);
}
else {
/* no look-ahead as this is last iteration */
trace_gpu_start( 0, 2, "syr2k", "syr2k last iteration" );
magma_ssyr2k(MagmaLower, MagmaNoTrans, pk, pk, c_neg_one,
da_ref(indi, indj), ldda,
dW , pm, d_one,
da_ref(indi, indi), ldda);
trace_gpu_end( 0, 2 );
}
indi_old = indi;
indj_old = indj;
pm_old = pm;
pn_old = pn;
} // end loop for(i)
/* Send the last block to the CPU */
pk = min(pm,pn);
if (1 <= n-nb){
spanel_to_q(MagmaUpper, pk-1, a_ref(n-pk+1, n-pk+2), lda, work);
trace_gpu_start( 0, 2, "get", "get last block" );
magma_sgetmatrix( pk, pk,
da_ref(n-pk+1, n-pk+1), ldda,
a_ref(n-pk+1, n-pk+1), lda );
trace_gpu_end( 0, 2 );
sq_to_panel(MagmaUpper, pk-1, a_ref(n-pk+1, n-pk+2), lda, work);
}
}// end of LOWER
trace_finalize( "ssytrd_sy2sb.svg", "trace.css" );
cudaEventDestroy(Pupdate_event);
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free( da );
MAGMA_S_SET2REAL( work[0], lwkopt );
magmablasSetKernelStream( 0 );
#if defined(USEMKL)
mkl_set_num_threads(1);
#endif
#if defined(USEACML)
omp_set_num_threads(1);
#endif
return *info;
} /* ssytrd_sy2sb_ */
/* ////////////////////////////////////////////////////////////////////////////
-- Testing magma_zhemm_mgpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
magmaDoubleComplex c_neg_one = MAGMA_Z_NEG_ONE;
magmaDoubleComplex calpha = MAGMA_Z_MAKE( 3.456, 5.678 );
magmaDoubleComplex cbeta = MAGMA_Z_MAKE( 1.234, 2.456 );
real_Double_t gflops, gpu_perf=0., cpu_perf=0., gpu_time=0., cpu_time=0.;
real_Double_t gpu_perf2=0., gpu_time2=0.;
double error=0., errorbis=0., work[1];
magmaDoubleComplex *hA, *hX, *hB, *hR;
magmaDoubleComplex_ptr dA[MagmaMaxGPUs], dX[MagmaMaxGPUs], dB[MagmaMaxGPUs], dwork[MagmaMaxGPUs], hwork[MagmaMaxGPUs+1];
magmaDoubleComplex_ptr dA2;
magma_int_t M, N, size, lda, ldda, msize, nb, nstream;
magma_int_t ione = 1;
magma_int_t iseed[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
parse_opts( argc, argv, &opts );
double tol = opts.tolerance * lapackf77_dlamch("E");
// default values
nb = (opts.nb > 0 ? opts.nb : 64);
nstream = (opts.nstream > 0 ? opts.nstream : 2);
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2];
magma_int_t nbcmplx = 0;
magma_buildconnection_mgpu(gnode, &nbcmplx, opts.ngpu);
printf("Initializing communication pattern... GPU-ncmplx %d\n\n", (int) nbcmplx);
for (int i=0; i < nbcmplx; ++i) {
int myngpu = gnode[i][MagmaMaxGPUs];
printf("cmplx %d has %d gpu ", i, myngpu);
for(int j=0; j < myngpu; ++j)
printf(" %d", (int) gnode[i][j]);
printf("\n");
}
magma_int_t nbevents = 2;
magma_queue_t streams[MagmaMaxGPUs][20];
magma_event_t redevents[MagmaMaxGPUs][20];
magma_event_t redevents2[MagmaMaxGPUs][MagmaMaxGPUs*MagmaMaxGPUs+10];
for( int d = 0; d < opts.ngpu; ++d ) {
for( magma_int_t i = 0; i < nstream; ++i ) {
magma_queue_create( &streams[d][i] );
}
for( magma_int_t i = 0; i < nbevents; ++i ) {
cudaEventCreateWithFlags(&redevents[d][i], cudaEventDisableTiming);
cudaEventCreateWithFlags(&redevents2[d][i], cudaEventDisableTiming);
}
}
printf( "nb %d, ngpu %d, nstream %d version %d\n", (int) nb, (int) opts.ngpu, (int) nstream, (int) opts.version );
printf(" M N nb offset CPU GFlop/s (sec) GPU GFlop/s (sec) CUBLAS hemm (sec) ||R|| / ||A||*||X||\n");
printf("=========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
M = opts.msize[itest];
N = opts.nsize[itest];
for( int offset = 0; offset < N; offset += min(N,nb) ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
msize = M - offset;
lda = M;
ldda = ((M + 31)/32)*32;
size = lda*M;
gflops = FLOPS_ZHEMM( MagmaLeft, (double)msize, (double)N ) / 1e9;
magma_int_t dworksiz = ldda*N*3;
magma_int_t hworksiz = lda*N;
TESTING_MALLOC_CPU( hA, magmaDoubleComplex, lda*M );
TESTING_MALLOC_CPU( hX, magmaDoubleComplex, lda*N );
TESTING_MALLOC_CPU( hB, magmaDoubleComplex, lda*N );
TESTING_MALLOC_PIN( hR, magmaDoubleComplex, lda*N );
for( int d = 0; d < opts.ngpu; ++d ) {
magma_int_t mlocal = ((M / nb) / opts.ngpu + 1) * nb;
magma_setdevice( d );
TESTING_MALLOC_DEV( dA[d], magmaDoubleComplex, ldda*mlocal );
TESTING_MALLOC_DEV( dX[d], magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( dB[d], magmaDoubleComplex, ldda*N );
TESTING_MALLOC_DEV( dwork[d], magmaDoubleComplex, dworksiz );
TESTING_MALLOC_PIN( hwork[d], magmaDoubleComplex, hworksiz );
}
TESTING_MALLOC_PIN( hwork[opts.ngpu], magmaDoubleComplex, lda*N );
if ( opts.check ) {
magma_setdevice( 0 );
TESTING_MALLOC_DEV( dA2, magmaDoubleComplex, ldda*M );
}
lapackf77_zlarnv( &ione, iseed, &size, hA );
magma_zmake_hermitian( M, hA, lda );
size = lda*N;
lapackf77_zlarnv( &ione, iseed, &size, hX );
lapackf77_zlarnv( &ione, iseed, &size, hB );
lapackf77_zlacpy( "Full", &M, &N, hB, &lda, hR, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_zsetmatrix_1D_col_bcyclic( M, M, hA, lda, dA, ldda, opts.ngpu, nb );
for( int d = 0; d < opts.ngpu; ++d ) {
magma_setdevice( d );
//magmablasSetKernelStream( streams[ d ][ 0 ] );
magma_zsetmatrix( M, N, hX, lda, dX[d], ldda );
//if (d == 0) magma_zsetmatrix( M, N, hB, lda, dB[d], ldda ); // this is wrong coz when offset != 0 the gpu who do the beta*C may be not 0 so this should be related to stdev(starting device who own i=0 first col)
magma_zsetmatrix( M, N, hB, lda, dB[d], ldda );
}
//memset(hR, 0, lda*N*sizeof(magmaDoubleComplex));
trace_init( 1, opts.ngpu, nstream, (magma_queue_t*) streams );
//magma_int_t offset = 0; //nb;
gpu_time = magma_sync_wtime(0);
magmablas_zhemm_mgpu_com(
MagmaLeft, MagmaLower, msize, N,
calpha, dA, ldda, offset,
dX, ldda,
cbeta, dB, ldda, dwork, dworksiz, hR, lda, hwork, hworksiz,
opts.ngpu, nb, streams, nstream, redevents2, nbevents, gnode, nbcmplx);
gpu_time = magma_sync_wtime(0) - gpu_time;
gpu_perf = gflops / gpu_time;
#ifdef TRACING
char buf[80];
snprintf( buf, sizeof(buf), "zhemm-m%d-n%d-nb%d-stream%d-ngpu%d-run%d.svg",
(int) M, (int) N, (int) nb, (int) nstream, (int) opts.ngpu, (int) iter );
trace_finalize( buf, "trace.css" );
#endif
/* ====================================================================
Performs operation using CUBLAS
=================================================================== */
if ( opts.check && iter == 0 ) {
magma_setdevice( 0 );
magmablasSetKernelStream( 0 );
magma_zsetmatrix( M, M, hA, lda, dA2, ldda );
magma_zsetmatrix( M, N, hX, lda, dX[0], ldda );
magma_zsetmatrix( M, N, hB, lda, dwork[0], ldda );
gpu_time2 = magma_sync_wtime(0);
magma_zhemm(
MagmaLeft, MagmaLower, msize, N,
calpha, dA2+offset*ldda+offset, ldda,
dX[0], ldda,
cbeta, dwork[0], ldda );
gpu_time2 = magma_sync_wtime(0) - gpu_time2;
gpu_perf2 = gflops / gpu_time2;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.check ) {
// store ||A||*||X||
errorbis = lapackf77_zlange("fro", &msize, &msize, hA+offset*lda+offset, &lda, work );
errorbis *= lapackf77_zlange("fro", &msize, &N, hX, &lda, work );
//printf( "A =" ); magma_zprint( M, M, hA, lda );
//printf( "X =" ); magma_zprint( M, N, hX, lda );
//printf( "B =" ); magma_zprint( M, N, hB, lda );
cpu_time = magma_wtime();
blasf77_zhemm( "Left", "Lower", &msize, &N,
&calpha, hA+offset*lda+offset, &lda,
hX, &lda,
&cbeta, hB, &lda );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
/*
trace_file = fopen("AJETE/C", "w");
for (int j = 0; j < N; j++)
for (int i = 0; i < siz; i++)
fprintf(trace_file, "%10d%10d%40.30e\n", i+1, j+1, hB[j*lda+i]);
fclose(trace_file);
*/
magma_int_t firstprint=0;
for(magma_int_t dev=0; dev < opts.ngpu; ++dev) {
magma_setdevice( dev );
magma_zgetmatrix( M, N, dB[dev], ldda, hR, lda );
// compute relative error ||R||/||A||*||X||, where R := B_magma - B_lapack = R - B
size = lda*N;
blasf77_zaxpy( &size, &c_neg_one, hB, &ione, hR, &ione );
error = lapackf77_zlange("fro", &msize, &N, hR, &lda, work) / errorbis;
//printf( "R =" ); magma_zprint( M, N, hR, lda );
if (firstprint == 0) {
printf( "%5d %5d %5d %5d %7.1f (%7.4f) %7.1f (%7.4f) %7.1f (%7.4f) %8.2e %s\n",
(int) M, (int) N, (int) nb, (int) offset,
cpu_perf, cpu_time,
gpu_perf, gpu_time,
gpu_perf2, gpu_time2,
error, (error < tol ? "ok" : "failed") );
}
else {
printf( "%89s %8.2e %s\n", " ",
error, (error < tol ? "ok" : "failed") );
}
status += ! (error < tol);
firstprint =1;
}
} else {
printf( "%5d %5d %5d %5d --- ( --- ) %7.1f (%7.4f) --- ( --- ) ---\n",
(int) M, (int) N, (int) nb, (int) offset,
gpu_perf, gpu_time );
}
TESTING_FREE_CPU( hA );
TESTING_FREE_CPU( hX );
TESTING_FREE_CPU( hB );
TESTING_FREE_PIN( hR );
for( int d = 0; d < opts.ngpu; ++d ) {
magma_setdevice( d );
TESTING_FREE_DEV( dA[d] );
TESTING_FREE_DEV( dX[d] );
TESTING_FREE_DEV( dB[d] );
TESTING_FREE_DEV( dwork[d] );
TESTING_FREE_PIN( hwork[d] );
}
TESTING_FREE_PIN( hwork[opts.ngpu] );
if ( opts.check ) {
magma_setdevice( 0 );
TESTING_FREE_DEV( dA2 );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
} // offset
printf( "\n" );
}
for( int d = 0; d < opts.ngpu; ++d ) {
magma_setdevice( d );
for( magma_int_t i = 0; i < nstream; ++i ) {
magma_queue_destroy( streams[d][i] );
}
for( magma_int_t i = 0; i < nbevents; ++i ) {
magma_event_destroy( redevents[d][i] );
magma_event_destroy( redevents2[d][i] );
}
}
TESTING_FINALIZE();
return status;
}
cudaError_t WINAPI wine_cudaEventCreateWithFlags( cudaEvent_t *event, int flags ) {
WINE_TRACE("\n");
return cudaEventCreateWithFlags( event, flags );
}
gaspi_return_t
gaspi_gpu_init(void)
{
gaspi_context_t * const gctx = &glb_gaspi_ctx;
int deviceCount;
cudaError_t cuda_error_id = cudaGetDeviceCount(&deviceCount);
if( cuda_error_id != cudaSuccess )
{
gaspi_print_error("Failed cudaGetDeviceCount." );
return GASPI_ERR_DEVICE;
}
if( deviceCount <= 0 )
{
gaspi_print_error("No CUDA capable devices found.");
return GASPI_ERR_DEVICE;
}
const int ib_numa_node = _gaspi_find_dev_numa_node();
int device_id = 0;
int gaspi_devices = 0;
int direct_devices[GPI2_GPU_MAX_DIRECT_DEVS];
struct cudaDeviceProp deviceProp;
for(device_id = 0; device_id < deviceCount; device_id++)
{
//TODO: possibly add functionality to show properties structure
cuda_error_id = cudaGetDeviceProperties(&deviceProp, device_id);
if( cuda_error_id != cudaSuccess)
{
return GASPI_ERR_DEVICE;
}
if( deviceProp.major >= 3 ) /* TODO: magic number */
{
cuda_error_id = cudaSetDevice(device_id);
if( cuda_error_id != cudaSuccess )
{
return GASPI_ERR_DEVICE;
}
if( ib_numa_node == _gaspi_find_GPU_numa_node(device_id) )
{
if( gaspi_devices < GPI2_GPU_MAX_DIRECT_DEVS - 1 )
{
direct_devices[gaspi_devices] = device_id;
gaspi_devices++;
}
}
}
}
if( 0 == gaspi_devices )
{
gaspi_print_error("No GPU Direct RDMA capable devices on the correct NUMA-socket were found.");
return GASPI_ERROR;
}
gpus = (gaspi_gpu_t*) malloc(sizeof(gaspi_gpu_t) * gaspi_devices);
if( gpus == NULL )
{
gaspi_print_error("Failed to allocate memory.");
return GASPI_ERR_MEMALLOC;
}
int i, j, k;
for(k = 0 ; k < gaspi_devices; k++)
{
cuda_error_id = cudaSetDevice(direct_devices[k]);
if( cuda_error_id != cudaSuccess )
{
return GASPI_ERR_DEVICE;
}
for(i = 0; i < GASPI_MAX_QP; i++)
{
cuda_error_id = cudaStreamCreate(&gpus[k].streams[i]);
if( cuda_error_id != cudaSuccess )
{
return GASPI_ERR_DEVICE;
}
for(j = 0; j < GASPI_CUDA_EVENTS; j++)
{
cuda_error_id = cudaEventCreateWithFlags(&gpus[k].events[i][j].event, cudaEventDisableTiming);
if( cuda_error_id != cudaSuccess )
{
return GASPI_ERR_DEVICE;
}
}
cuda_error_id = cudaStreamCreateWithFlags(&gpus[k].streams[i], cudaStreamNonBlocking);
if( cuda_error_id != cudaSuccess )
{
return GASPI_ERR_DEVICE;
}
}
gpus[k].device_id = direct_devices[k];
}
gctx->gpu_count = gaspi_devices;
gctx->use_gpus = 1;
return GASPI_SUCCESS;
}
extern "C" magma_int_t
magma_sbulge_applyQ_v2(char side,
magma_int_t NE, magma_int_t N,
magma_int_t NB, magma_int_t Vblksiz,
float *dE, magma_int_t ldde,
float *V, magma_int_t ldv,
float *T, magma_int_t ldt,
magma_int_t *info)
{
//%===========================
//% local variables
//%===========================
magma_int_t Vm, Vn, mt, nt;
magma_int_t myrow, mycol, blkj, blki;
magma_int_t blkid,vpos,tpos;
magma_int_t firstrow, nbcolinvolvd;
magma_int_t versionL = 113;
magma_int_t versionR = 92;
magma_int_t Vchunksiz = 10;
*info=0;
/* Quick return */
if ( NE == 0 ) {
return MAGMA_SUCCESS;
}
if ( N == 0 ) {
return MAGMA_SUCCESS;
}
if ( NB == 0 ) {
return MAGMA_SUCCESS;
}
/* ==========================================
* some infos for developer
* Initialisation and checking nb of cores
* ==========================================*/
/* we have 2 algo for left (113 114) and 2 algo for right (91 92)
* which correspond to versionL versionR.
* They are very similar (detail explained in tech report and matlab code)
* however version 114 and 92 improve locality.
* while version 113 is used in case WNATZ=1 (construct Q2) which allow
* the construction to be done in an optimized way taking into
* consideration that the matrix is Identity so making less flops.
*
*/
// Initialize streaming and events
cudaDeviceSynchronize();
magma_queue_t cstream;
magmablasGetKernelStream(&cstream);
magma_queue_t stream[2];
magma_queue_create( &stream[0] );
magma_queue_create( &stream[1] );
magma_event_t myevent[2];
cudaEventCreateWithFlags(&myevent[0],cudaEventDisableTiming);
cudaEventCreateWithFlags(&myevent[1],cudaEventDisableTiming);
// Azzam 21/11/2012
// NOTE THAT dwork was of size 2*NE*Vblksiz+...
// but I am thinking why not modifing it to NE*Vblksiz+...
// BUT NO because the 2* is used because of making 2 streams working and so
// they might be using dwork in parallel
float *dwork, *dwork0, *dwork1, *dwvt0, *dwvt1;
float *dT0, *dV0, *dT1, *dV1;
magma_int_t lddv = ldv;
magma_int_t lddt = ldt;
magma_int_t lddw = 0;
magma_int_t lddwork = ((NE+31)/32)*32;
magma_int_t dwVTsiz = lddv*Vblksiz; // lddv*lddv + lddv*lddwork;(v2) // lddv*Vblksiz; (v1,v3)
magma_int_t dworksiz = lddwork*Vblksiz; // lddv*Vblksiz; (v2) // NE*Vblksiz=lddwork*Vblksiz; (v1,v3)
if(MAGMA_SUCCESS != magma_smalloc( &dwork, 2*dworksiz + 2*dwVTsiz + 2*Vchunksiz* (Vblksiz* (lddv+lddt)) )) {
printf ("!!!! magma_sbulge_applyQ magma_alloc failed for: dwork\n" );
exit(-1);
}
dwork0 = dwork; // size = dworksiz;
dwork1 = dwork0 + dworksiz; // size = dworksiz;
dwvt0 = dwork + 2*dworksiz; // size = dwVTsiz;
dwvt1 = dwvt0 + dwVTsiz; // size = dwVTsiz;
dV0 = dwork + 2*dworksiz + 2*dwVTsiz;
dT0 = dV0 + Vchunksiz*Vblksiz*lddv;
dV1 = dT0 + Vchunksiz*Vblksiz*lddt;
dT1 = dV1 + Vchunksiz*Vblksiz*lddv;
// make overlapped copy
magma_int_t ncpy = 0;
magma_int_t copyed=0, copyst=0;
magma_int_t blkcnt,nothing, mysiz, flip, vld,tld, locpos;
findVTsiz(N, NB, Vblksiz, &blkcnt, ¬hing);
flip = 0;
// performance loss if the reflector are applied to a big number of eigenvectors (~10000)
// => apply the reflectors to blocks of eigenvectors.
//magma_int_t nr_bl = magma_ceildiv(NE,10000); //nr of blocks
magma_int_t sz_bl = NE; //magma_ceildiv(NE,nr_bl*64)*64; //maximum size of blocks (to have blocks of around the same size and multiple of 64)
magma_int_t ib; //size of current block
/* SIDE LEFT meaning apply E = Q*E = (q_1*q_2*.....*q_n) * E ==> so traverse Vs in reverse order (forward) from q_n to q_1
* Also E is splitten by row meaning each apply consist in a block of row (horizontal block) */
/* SIDE RIGHT meaning apply E = E*Q = E * (q_1*q_2*.....*q_n) ==> so tarverse Vs in normal order (forward) from q_1 to q_n
* Also E is splitten by col meaning each apply consist in a block of col (vertical block) */
#ifdef ENABLE_DEBUG
printf(" APPLY Q_v22 GPU with N %d, NE %d, NB %d, Vblksiz %d, versionL %d versionR %d SIDE %c \n",
N, NE, NB, Vblksiz, versionL, versionR, side);
#endif
/*
* MagmamaLeft
*/
if(side=='L'){
/*
* Version 113:
* loop over the block_col (nt) and for each find the
* number of tiles (mt) in this block_col. then loop over mt, find
* the size of the V's(Vm,Vn) and apply it to the corresponding
* portion of E.
*/
if( versionL == 113 ) {
nt = magma_ceildiv((N-1),Vblksiz);
for (blkj=nt-1; blkj>=0; blkj--) {
/* the index of the first row on the top of block (blkj) */
firstrow = blkj * Vblksiz + 1;
/*find the number of tile for this block */
if( blkj == nt-1 )
mt = magma_ceildiv( N - firstrow, NB);
else
mt = magma_ceildiv( N - (firstrow+1), NB);
/*loop over the tiles find the size of the Vs and apply it */
for (blki=mt; blki>0; blki--) {
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + (mt-blki)*NB;
mycol = blkj*Vblksiz;
Vm = min( NB+Vblksiz-1, N-myrow);
if( ( blkj == nt-1 ) && ( blki == mt ) ){
Vn = min (Vblksiz, Vm);
} else {
Vn = min (Vblksiz, Vm-1);
}
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
//printf("voici blkj %d blki %d Vm %d Vn %d mycol %d vpos %d \n",blkj,blki,Vm, Vn,mycol,vpos);
magma_bulge_findpos113(N, NB, Vblksiz, mycol, myrow, &blkid);
// COPY Vchunksiz Vs and Vchunksiz Ts to GPU and store it in dV0/dV1 and dT0/dT1
if(ncpy==0){
// flip = 1 for this.
copyst = 0; // meaning that copy will start copying from blkid =copyst
copyed = min(copyst+Vchunksiz, blkcnt); // meaning that copy will end copying at blkid =copyed-1==> next copy had to start at copyed
mysiz = copyed-copyst; // the size of the chunk to be copied
if(mysiz>0){
ncpy = 1;
flip = 1;
vpos = copyst*Vblksiz*ldv;
tpos = copyst*Vblksiz*ldt;
vld = mysiz * ldv;
tld = mysiz * ldt;
magmablasSetKernelStream(stream[1]);
magma_ssetmatrix_async(vld, Vblksiz, V(vpos), vld, dV1, vld, stream[1]);
magma_ssetmatrix_async(tld, Vblksiz, T(tpos), tld, dT1, tld, stream[1]);
//printf("doing the first copy of mysiz %2d copyst %2d copyed %2d vpos %8d tpos %8d into dV1 dT1\n",mysiz,copyst,copyed,vpos,tpos);
}
}
if(blkid == copyst){
flip = ncpy % 2;
copyst = copyed; // meaning that copy will start copying from blkid =copyst
copyed = min(copyst+Vchunksiz, blkcnt); // meaning that copy will end copying at blkid =copyed-1==> next copy had to start at copyed
mysiz = copyed-copyst; // the size of the chunk to be copied
//printf(" get to copy blkid %d blkid+(2*Vchunksiz) %d copyst %d copyed %d\n",blkid,blkid+(Vchunksiz),copyst,copyed);
if(mysiz>0){
ncpy = ncpy + 1;
vpos = copyst*Vblksiz*ldv;
tpos = copyst*Vblksiz*ldt;
vld = mysiz * ldv;
tld = mysiz * ldt;
if(flip==0){ // now I am working on dV0 so copy the next and put it on dV1
//printf("doing overlapping copy of mysiz %2d copyst %2d copyed %2d vpos %8d tpos %8d into dV1 dT1\n",mysiz,copyst,copyed,vpos,tpos);
magmablasSetKernelStream(stream[1]);
magma_ssetmatrix_async(vld, Vblksiz, V(vpos), vld, dV1, vld, stream[1]);
magma_ssetmatrix_async(tld, Vblksiz, T(tpos), tld, dT1, tld, stream[1]);
}else{ // now I am working on dV1 so copy the next and put it on dV0
//printf("doing overlapping copy of mysiz %2d copyst %2d copyed %2d vpos %8d tpos %8d into dV0 dT0\n",mysiz,copyst,copyed,vpos,tpos);
magmablasSetKernelStream(stream[0]);
magma_ssetmatrix_async(vld, Vblksiz, V(vpos), vld, dV0, vld, stream[0]);
magma_ssetmatrix_async(tld, Vblksiz, T(tpos), tld, dT0, tld, stream[0]);
}
}
}
if((Vm>0)&&(Vn>0)){
locpos = blkid%Vchunksiz;
magma_int_t lcvpos = locpos*Vblksiz*lddv;
magma_int_t lctpos = locpos*Vblksiz*lddt;
//printf("voici blkj %d blki %d Vm %d Vn %d mycol %d locvpos %5d loctpos %5d blkid %2d using data in dV%1d dT%1d \n",blkj,blki,Vm, Vn,mycol,lcvpos,lctpos, blkid,flip,flip);
if(flip==0){
magmablasSetKernelStream(stream[0]);
cudaStreamWaitEvent(stream[0], myevent[1], 0);
for(magma_int_t i=0; i<NE; i+= sz_bl){
ib = min(sz_bl, NE-i);
lddw = min(lddwork,sz_bl);
//magma_slarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV0+lcvpos, lddv, dT0+lctpos, lddt, dE(myrow,i), ldde, dwork0, lddw);
magma_slarfb_gpu_gemm( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV0+lcvpos, lddv, dT0+lctpos, lddt, dE(myrow,i), ldde, dwork0, lddw, dwvt0, lddv);
}
cudaEventRecord(myevent[0], stream[0]);
}else{
magmablasSetKernelStream(stream[1]);
cudaStreamWaitEvent(stream[1], myevent[0], 0);
for(magma_int_t i=0; i<NE; i+= sz_bl){
ib = min(sz_bl, NE-i);
lddw = min(lddwork,sz_bl);
//magma_slarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV1+lcvpos, lddv, dT1+lctpos, lddt, dE(myrow,i), ldde, dwork1, lddw);
magma_slarfb_gpu_gemm( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV1+lcvpos, lddv, dT1+lctpos, lddt, dE(myrow,i), ldde, dwork1, lddw, dwvt1, lddv);
}
cudaEventRecord(myevent[1], stream[1]);
}
} // end for (Vm &Vn) > 0
} // end for blki
} // end for blkj
} // end if version=113
/*
* Version 114:
* loop over the block_row (mt) and for each find diagonally the
* number of tiles (nt) in this block_row. then loop over nt, find
* the size of the V's(Vm,Vn) and apply it to the corresponding
* portion of E.
*/
else {
mt = magma_ceildiv((N-1),NB);
for (blki = mt; blki>0; blki--) {
/* nbcolinvolvd = number of column corresponding to this block_row (blki) */
nbcolinvolvd = min(N-1, blki*NB);
/*find the number of tile for this block (diagonal row of tiles) */
nt = magma_ceildiv(nbcolinvolvd,Vblksiz);
/*loop over the tiles find the size of the Vs and apply it */
for (blkj = nt-1; blkj>=0; blkj--) {
/* the index of the first row of the first col meaning
* the block on the top left (blki) */
firstrow = (mt-blki)*NB+1;
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + blkj*Vblksiz;
mycol = blkj*Vblksiz;
Vm = min( NB+Vblksiz-1, N-myrow);
if( ( blkj == nt-1 ) && ( blki == mt ) ){
Vn = min (Vblksiz, Vm);
}else{
Vn = min (Vblksiz, Vm-1);
}
if((Vm>0)&&(Vn>0)){
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
magma_bulge_findVTpos(N, NB, Vblksiz, mycol, myrow, ldv, ldt, &vpos, &tpos);
magma_ssetmatrix_async(Vm, Vn, V(vpos), ldv, dV0, lddv, NULL);
magma_ssetmatrix_async(Vn, Vn, T(tpos), ldt, dT0, lddt, NULL);
//printf("voici blki %d rownbm %d mycol %d coled %d blkid %d vpos %d tpos %d\n", blki, rownbm, mycol, coled, blkid, vpos, tpos);
for(magma_int_t i=0; i<NE; i+= sz_bl){
ib = min(sz_bl, NE-i);
magma_slarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV0, lddv, dT0, lddt, dE(myrow,i), ldde, dwork, NE);
}
} // end for (Vm &Vn) > 0
} // end for blkj
} // end for blki
} // end version 114
} // end LEFT
/*
* MagmaRight
*/
else {
/*
* Version 91:
*/
if( versionR == 91 ) {
nt = magma_ceildiv((N-1),Vblksiz);
for (blkj=0; blkj<nt; blkj++) {
/* the index of the first myrow on the top of block (blkj) */
firstrow = blkj * Vblksiz + 1;
/*find the number of tile for this block */
if( blkj == nt-1 )
mt = magma_ceildiv( N - firstrow, NB);
else
mt = magma_ceildiv( N - (firstrow+1), NB);
/*loop over the tiles find the size of the Vs and apply it */
for (blki=1; blki<=mt; blki++) {
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + (mt-blki)*NB;
Vm = min( NB+Vblksiz-1, N-myrow);
if( ( blkj == nt-1 ) && ( blki == mt ) )
{
Vn = min (Vblksiz, Vm);
}else{
Vn = min (Vblksiz, Vm-1);
}
mycol = blkj*Vblksiz;
if((Vm>0)&&(Vn>0)){
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
magma_bulge_findVTpos(N, NB, Vblksiz ,mycol, myrow, ldv, ldt, &vpos, &tpos);
magma_ssetmatrix_async(Vm, Vn, V(vpos), ldv, dV0, lddv, NULL);
magma_ssetmatrix_async(Vn, Vn, T(tpos), ldt, dT0, lddt, NULL);
magma_slarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, NE, Vm, Vn, dV0, lddv, dT0, lddt, dE(0, myrow), ldde, dwork, NE);
} // end for (Vm &Vn) > 0
} // end for blki
} // end fo blkj
} // end of version 91
/*
* Version 92:
*/
else {
mt = magma_ceildiv((N-1),NB);
for (blki = 1; blki<=mt; blki++) {
/* nbcolinvolvd = number of column corresponding to this block_row (blki) */
nbcolinvolvd = min(N-1, blki*NB);
/*find the number of tile for this block (diagonal row of tiles) */
nt = magma_ceildiv(nbcolinvolvd,Vblksiz);
/*loop over the tiles find the size of the Vs and apply it */
for (blkj = 0; blkj<nt; blkj++) {
/* the index of the first row of the first col meaning
* the block on the top left (blki) */
firstrow = (mt-blki)*NB+1;
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + blkj*Vblksiz;
mycol = blkj*Vblksiz;
Vm = min( NB+Vblksiz-1, N-myrow);
if( ( blkj == nt-1 ) && ( blki == mt ) ){
Vn = min (Vblksiz, Vm);
}else{
Vn = min (Vblksiz, Vm-1);
}
if((Vm>0)&&(Vn>0)){
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
magma_bulge_findVTpos(N, NB, Vblksiz ,mycol, myrow, ldv, ldt, &vpos, &tpos);
magma_ssetmatrix_async(Vm, Vn, V(vpos), ldv, dV0, lddv, NULL);
magma_ssetmatrix_async(Vn, Vn, T(tpos), ldt, dT0, lddt, NULL);
magma_slarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, NE, Vm, Vn, dV0, lddv, dT0, lddt, dE(0, myrow), ldde, dwork, NE);
} // end for (Vm &Vn) > 0
} //end for blkj
} // end for blki
} //end of version 92
} // end RIGHT
cudaDeviceSynchronize();
magmablasSetKernelStream(cstream);
cudaEventDestroy(myevent[0]);
cudaEventDestroy(myevent[1]);
magma_queue_destroy( stream[0] );
magma_queue_destroy( stream[1] );
magma_free(dwork);
return MAGMA_SUCCESS;
}
static void _XMP_reflect_pcopy_sched_dim(_XMP_array_t *adesc, int target_dim,
int lwidth, int uwidth, int is_periodic, void *dev_array_addr, int *lwidths, int *uwidths){
//printf("desc=%p, tardim=%d, lw=%d, uw=%d, devp=%p\n", adesc, target_dim, lwidth, uwidth, dev_array_addr);
if (lwidth == 0 && uwidth == 0) return;
_XMP_array_info_t *ai = &(adesc->info[target_dim]);
_XMP_array_info_t *ainfo = adesc->info;
_XMP_ASSERT(ai->align_manner == _XMP_N_ALIGN_BLOCK);
_XMP_ASSERT(ai->is_shadow_comm_member);
if (lwidth > ai->shadow_size_lo || uwidth > ai->shadow_size_hi){
_XMP_fatal("reflect width is larger than shadow width.");
}
_XMP_reflect_sched_t *reflect = ai->reflect_acc_sched;
int target_tdim = ai->align_template_index;
_XMP_nodes_info_t *ni = adesc->align_template->chunk[target_tdim].onto_nodes_info;
int ndims = adesc->dim;
// 0-origin
int my_pos = ni->rank;
int lb_pos = _XMP_get_owner_pos(adesc, target_dim, ai->ser_lower);
int ub_pos = _XMP_get_owner_pos(adesc, target_dim, ai->ser_upper);
int lo_pos = (my_pos == lb_pos) ? ub_pos : my_pos - 1;
int hi_pos = (my_pos == ub_pos) ? lb_pos : my_pos + 1;
MPI_Comm *comm = adesc->align_template->onto_nodes->comm;
int my_rank = adesc->align_template->onto_nodes->comm_rank;
int lo_rank = my_rank + (lo_pos - my_pos) * ni->multiplier;
int hi_rank = my_rank + (hi_pos - my_pos) * ni->multiplier;
int type_size = adesc->type_size;
//void *array_addr = adesc->array_addr_p;
void *lo_send_array = NULL;
void *lo_recv_array = NULL;
void *hi_send_array = NULL;
void *hi_recv_array = NULL;
void *lo_send_dev_buf = NULL;
void *lo_recv_dev_buf = NULL;
void *hi_send_dev_buf = NULL;
void *hi_recv_dev_buf = NULL;
void *lo_send_host_buf = NULL;
void *lo_recv_host_buf = NULL;
void *hi_send_host_buf = NULL;
void *hi_recv_host_buf = NULL;
void *mpi_lo_send_buf = NULL;
void *mpi_lo_recv_buf = NULL;
void *mpi_hi_send_buf = NULL;
void *mpi_hi_recv_buf = NULL;
int lo_buf_size = 0;
int hi_buf_size = 0;
//
// setup data_type
//
int count = 0, blocklength = 0;
long long stride = 0;
// int count_offset = 0;
if (_XMPF_running && !_XMPC_running){ /* for XMP/F */
count = 1;
blocklength = type_size;
stride = ainfo[0].alloc_size * type_size;
for (int i = ndims - 2; i >= target_dim; i--){
count *= ainfo[i+1].alloc_size;
}
for (int i = 1; i <= target_dim; i++){
blocklength *= ainfo[i-1].alloc_size;
stride *= ainfo[i].alloc_size;
}
}
else if (!_XMPF_running && _XMPC_running){ /* for XMP/C */
count = 1;
blocklength = type_size;
stride = ainfo[ndims-1].alloc_size * type_size;
/* if(target_dim > 0){ */
/* count *= ainfo[0].par_size; */
/* count_offset = ainfo[0].shadow_size_lo; */
/* } */
/* for (int i = 1; i < target_dim; i++){ */
/* count *= ainfo[i].alloc_size; */
/* } */
/* for (int i = ndims - 2; i >= target_dim; i--){ */
/* blocklength *= ainfo[i+1].alloc_size; */
/* stride *= ainfo[i].alloc_size; */
/* } */
if(target_dim == 0){
count *= 1;
if(ndims >= 2){
blocklength *= (ainfo[1].par_size + lwidths[1] + uwidths[1]);
}
}else{
count *= (ainfo[0].par_size + lwidths[0] + uwidths[0]);
for(int i = 1; i < target_dim; i++){
count *= ainfo[i].alloc_size;
}
blocklength *= ainfo[target_dim+1].alloc_size;
stride *= ainfo[target_dim].alloc_size;
}
for(int i = target_dim+2; i < ndims; i++){
blocklength *= ainfo[i].alloc_size;
}
for(int i = target_dim+1 ; i < ndims - 1; i++){
stride *= ainfo[i].alloc_size;
}
/* mod_4 */
count = 1;
blocklength = 1;
stride = 1;
for(int i = 0; i < ndims; i++){
int fact = (i == target_dim)? 1 : (ainfo[i].par_size + lwidths[i] + uwidths[i]);
int alloc_size = ainfo[i].alloc_size;
if(blocklength == 1 || fact == alloc_size){
blocklength *= fact;
stride *= alloc_size;
}else if(count == 1 && target_dim != 0){ //to be contiguous if target_dim==0
count = blocklength;
blocklength = fact;
stride = alloc_size;
}else{
blocklength *= alloc_size;
stride *= alloc_size;
}
//printf("tar=%d, i=%d, fact=%d, allocsize=%d, (%d,%d,%lld)\n", target_dim, i, fact, alloc_size, count , blocklength, stride);
}
blocklength *= type_size;
stride *= type_size;
/* mod_4 end */
/* it used at 150717
for (int i = 1; i <= target_dim; i++){
count *= ainfo[i-1].alloc_size;
}
for (int i = ndims - 2; i >= target_dim; i--){
blocklength *= ainfo[i+1].alloc_size;
stride *= ainfo[i].alloc_size;
}
*/
/* for (int i = target_dim + 1; i < ndims; i++){ */
/* blocklength *= ainfo[i].alloc_size; */
/* } */
/* for (int i = target_dim; i < ndims - 1; i++){ */
/* stride *= ainfo[i].alloc_size; */
/* } */
// printf("count =%d, blength=%d, stride=%lld\n", count ,blocklength, stride);
// printf("ainfo[0].par_size=%d\n", ainfo[0].par_size);
// printf("count_ofset=%d,\n", count_offset);
}
else {
_XMP_fatal("cannot determin the base language.");
}
//
// calculate base address
//
// for lower reflect
if (lwidth){
lo_send_array = lo_recv_array = (void *)((char*)dev_array_addr + /*count_offset*/0 * stride);
for (int i = 0; i < ndims; i++) {
int lb_send, lb_recv;
unsigned long long dim_acc;
if (i == target_dim) {
//printf("ainfo[%d].local_upper=%d\n",i,ainfo[i].local_upper);
lb_send = ainfo[i].local_upper - lwidth + 1;
lb_recv = ainfo[i].shadow_size_lo - lwidth; ////ainfo[i].local_lower - lwidth;
} else {
// Note: including shadow area
lb_send = 0; //// ainfo[i].local_lower - ainfo[i].shadow_size_lo;
lb_recv = 0; //// ainfo[i].local_lower - ainfo[i].shadow_size_lo;
}
dim_acc = ainfo[i].dim_acc;
lo_send_array = (void *)((char *)lo_send_array + lb_send * dim_acc * type_size);
lo_recv_array = (void *)((char *)lo_recv_array + lb_recv * dim_acc * type_size);
}
}
// for upper reflect
if (uwidth){
hi_send_array = hi_recv_array = (void *)((char*)dev_array_addr + /*count_offset*/0 * stride);
for (int i = 0; i < ndims; i++) {
int lb_send, lb_recv;
unsigned long long dim_acc;
if (i == target_dim) {
lb_send = ainfo[i].local_lower;
lb_recv = ainfo[i].local_upper + 1;
} else {
// Note: including shadow area
lb_send = 0; //ainfo[i].local_lower - ainfo[i].shadow_size_lo;
lb_recv = 0; //ainfo[i].local_lower - ainfo[i].shadow_size_lo;
}
dim_acc = ainfo[i].dim_acc;
hi_send_array = (void *)((char *)hi_send_array + lb_send * dim_acc * type_size);
hi_recv_array = (void *)((char *)hi_recv_array + lb_recv * dim_acc * type_size);
}
}
// for lower reflect
if (reflect->datatype_lo != MPI_DATATYPE_NULL){
MPI_Type_free(&reflect->datatype_lo);
}
if(packVector || count == 1){
MPI_Type_contiguous(blocklength * lwidth * count, MPI_BYTE, &reflect->datatype_lo);
// MPI_Type_contiguous(blocklength * lwidth * count / type_size, MPI_FLOAT, &reflect->datatype_lo);
fprintf(stderr, "dim=%d, send elements lo = %d\n", target_dim, blocklength * lwidth * count / type_size);
//fprintf(stderr, "useHostBuf=%c , packVector=%c\n", useHostBuffer, packVector);
// if(useHostBuffer){ fprintf(stderr,"using host buffer\n"); }
// if(packVector){ fprintf(stderr, "using pack vector\n"); }
}else{
MPI_Type_vector(count, blocklength * lwidth, stride, MPI_BYTE, &reflect->datatype_lo);
}
MPI_Type_commit(&reflect->datatype_lo);
// for upper reflect
if (reflect->datatype_hi != MPI_DATATYPE_NULL){
MPI_Type_free(&reflect->datatype_hi);
}
if(packVector || count == 1){
MPI_Type_contiguous(blocklength * uwidth * count, MPI_BYTE, &reflect->datatype_hi);
// MPI_Type_contiguous(blocklength * uwidth * count / type_size, MPI_FLOAT, &reflect->datatype_hi);
fprintf(stderr, "dim=%d, send elements hi = %d\n", target_dim, blocklength * uwidth * count / type_size);
}else{
MPI_Type_vector(count, blocklength * uwidth, stride, MPI_BYTE, &reflect->datatype_hi);
}
MPI_Type_commit(&reflect->datatype_hi);
//
// Allocate buffers
//
if(useHostBuffer){
CUDA_SAFE_CALL(cudaFreeHost(reflect->lo_send_host_buf));
CUDA_SAFE_CALL(cudaFreeHost(reflect->lo_recv_host_buf));
}
if ((_XMPF_running && target_dim != ndims - 1) ||
(_XMPC_running && target_dim != 0)){
if(packVector){
CUDA_SAFE_CALL(cudaFree(reflect->lo_send_buf));
CUDA_SAFE_CALL(cudaFree(reflect->lo_recv_buf));
}
}
if ((_XMPF_running && target_dim == ndims - 1) ||
(_XMPC_running && target_dim == 0)){
//
}
// for lower reflect
if (lwidth){
lo_buf_size = lwidth * blocklength * count;
hi_buf_size = uwidth * blocklength * count;
if ((_XMPF_running && target_dim == ndims - 1) ||
(_XMPC_running && target_dim == 0)){
lo_send_dev_buf = lo_send_array;
lo_recv_dev_buf = lo_recv_array;
hi_send_dev_buf = hi_send_array;
hi_recv_dev_buf = hi_recv_array;
} else {
_XMP_TSTART(t0);
if(packVector){
CUDA_SAFE_CALL(cudaMalloc((void **)&lo_send_dev_buf, lo_buf_size + hi_buf_size));
hi_send_dev_buf = (char*)lo_send_dev_buf + lo_buf_size;
CUDA_SAFE_CALL(cudaMalloc((void **)&lo_recv_dev_buf, lo_buf_size + hi_buf_size));
hi_recv_dev_buf = (char*)lo_recv_dev_buf + lo_buf_size;
}else{
lo_send_dev_buf = lo_send_array;
lo_recv_dev_buf = lo_recv_array;
hi_send_dev_buf = hi_send_array;
hi_recv_dev_buf = hi_recv_array;
}
_XMP_TEND2(xmptiming_.t_mem, xmptiming_.tdim_mem[target_dim], t0);
}
if(useHostBuffer){
CUDA_SAFE_CALL(cudaMallocHost((void**)&lo_send_host_buf, lo_buf_size + hi_buf_size));
hi_send_host_buf = (char*)lo_send_host_buf + lo_buf_size;
CUDA_SAFE_CALL(cudaMallocHost((void**)&lo_recv_host_buf, lo_buf_size + hi_buf_size));
hi_recv_host_buf = (char*)lo_recv_host_buf + lo_buf_size;
mpi_lo_send_buf = lo_send_host_buf;
mpi_lo_recv_buf = lo_recv_host_buf;
mpi_hi_send_buf = hi_send_host_buf;
mpi_hi_recv_buf = hi_recv_host_buf;
}else{
mpi_lo_send_buf = lo_send_dev_buf;
mpi_lo_recv_buf = lo_recv_dev_buf;
mpi_hi_send_buf = hi_send_dev_buf;
mpi_hi_recv_buf = hi_recv_dev_buf;
}
}
// for upper reflect
//
// initialize communication
//
int src, dst;
if (!is_periodic && my_pos == lb_pos){ // no periodic
lo_rank = MPI_PROC_NULL;
}
if (!is_periodic && my_pos == ub_pos){ // no periodic
hi_rank = MPI_PROC_NULL;
}
// for lower shadow
if (lwidth){
src = lo_rank;
dst = hi_rank;
} else {
src = MPI_PROC_NULL;
dst = MPI_PROC_NULL;
}
// fprintf(stderr, "dim=%d, lo_src=%d, lo_dst=%d\n", target_dim, src, dst);
if (reflect->req[0] != MPI_REQUEST_NULL){
MPI_Request_free(&reflect->req[0]);
}
if (reflect->req[1] != MPI_REQUEST_NULL){
MPI_Request_free(&reflect->req[1]);
}
MPI_Recv_init(mpi_lo_recv_buf, 1, reflect->datatype_lo, src,
_XMP_N_MPI_TAG_REFLECT_LO, *comm, &reflect->req[0]);
MPI_Send_init(mpi_lo_send_buf, 1, reflect->datatype_lo, dst,
_XMP_N_MPI_TAG_REFLECT_LO, *comm, &reflect->req[1]);
// for upper shadow
if (uwidth){
src = hi_rank;
dst = lo_rank;
} else {
src = MPI_PROC_NULL;
dst = MPI_PROC_NULL;
}
// fprintf(stderr, "dim=%d, hi_src=%d, hi_dst=%d\n", target_dim, src, dst);
if (reflect->req[2] != MPI_REQUEST_NULL){
MPI_Request_free(&reflect->req[2]);
}
if (reflect->req[3] != MPI_REQUEST_NULL){
MPI_Request_free(&reflect->req[3]);
}
MPI_Recv_init(mpi_hi_recv_buf, 1, reflect->datatype_hi, src,
_XMP_N_MPI_TAG_REFLECT_HI, *comm, &reflect->req[2]);
MPI_Send_init(mpi_hi_send_buf, 1, reflect->datatype_hi, dst,
_XMP_N_MPI_TAG_REFLECT_HI, *comm, &reflect->req[3]);
//
// cache schedule
//
reflect->count = count;
reflect->blocklength = blocklength;
reflect->stride = stride;
reflect->lo_send_array = lo_send_array;
reflect->lo_recv_array = lo_recv_array;
reflect->hi_send_array = hi_send_array;
reflect->hi_recv_array = hi_recv_array;
if(packVector){
reflect->lo_send_buf = lo_send_dev_buf;
reflect->lo_recv_buf = lo_recv_dev_buf;
reflect->hi_send_buf = hi_send_dev_buf;
reflect->hi_recv_buf = hi_recv_dev_buf;
}
if(useHostBuffer){
reflect->lo_send_host_buf = lo_send_host_buf;
reflect->lo_recv_host_buf = lo_recv_host_buf;
reflect->hi_send_host_buf = hi_send_host_buf;
reflect->hi_recv_host_buf = hi_recv_host_buf;
}
reflect->lo_rank = lo_rank;
reflect->hi_rank = hi_rank;
// gpu async
reflect->lo_async_id = _XMP_alloc(sizeof(cudaStream_t));
CUDA_SAFE_CALL(cudaStreamCreate(reflect->lo_async_id));
if(target_dim != 0 &&
(!useHostBuffer || (lo_rank != MPI_PROC_NULL && hi_rank != MPI_PROC_NULL && (lo_buf_size / type_size) <= useSingleStreamLimit)) ){
reflect->hi_async_id = NULL;
}else{
cudaStream_t *hi_stream = (cudaStream_t*)_XMP_alloc(sizeof(cudaStream_t));
CUDA_SAFE_CALL(cudaStreamCreate(hi_stream));
reflect->hi_async_id = (void*)hi_stream;
}
reflect->event = _XMP_alloc(sizeof(cudaEvent_t));
CUDA_SAFE_CALL(cudaEventCreateWithFlags(reflect->event, cudaEventDisableTiming));
}
CudaEvent::CudaEvent( ) : isRecorded( false ), finished( true ), refCounter( 0u )
{
log( ggLog::CUDA_RT()+ggLog::MEMORY(), "create event" );
CUDA_CHECK( cudaEventCreateWithFlags( &event, cudaEventDisableTiming ) );
}
static int copy_data_1_to_1_generic(starpu_data_handle_t handle,
struct _starpu_data_replicate *src_replicate,
struct _starpu_data_replicate *dst_replicate,
struct _starpu_data_request *req)
{
unsigned src_node = src_replicate->memory_node;
unsigned dst_node = dst_replicate->memory_node;
STARPU_ASSERT(src_replicate->refcnt);
STARPU_ASSERT(dst_replicate->refcnt);
STARPU_ASSERT(src_replicate->allocated);
STARPU_ASSERT(dst_replicate->allocated);
_starpu_comm_amounts_inc(src_node, dst_node, handle->ops->get_size(handle));
#ifdef STARPU_SIMGRID
return _starpu_simgrid_transfer(handle->ops->get_size(handle), src_node, dst_node, req);
#else /* !SIMGRID */
int ret = 0;
const struct starpu_data_copy_methods *copy_methods = handle->ops->copy_methods;
enum starpu_node_kind src_kind = starpu_node_get_kind(src_node);
enum starpu_node_kind dst_kind = starpu_node_get_kind(dst_node);
#ifdef STARPU_USE_CUDA
cudaError_t cures;
cudaStream_t stream;
#endif
void *src_interface = src_replicate->data_interface;
void *dst_interface = dst_replicate->data_interface;
#if defined(STARPU_USE_CUDA) && defined(HAVE_CUDA_MEMCPY_PEER) && !defined(STARPU_SIMGRID)
if ((src_kind == STARPU_CUDA_RAM) || (dst_kind == STARPU_CUDA_RAM))
{
unsigned devid;
if ((src_kind == STARPU_CUDA_RAM) && (dst_kind == STARPU_CUDA_RAM))
{
/* GPU-GPU transfer, issue it from the device we are supposed to drive */
int worker = starpu_worker_get_id();
devid = starpu_worker_get_devid(worker);
}
else
{
unsigned node = (dst_kind == STARPU_CUDA_RAM)?dst_node:src_node;
devid = _starpu_memory_node_get_devid(node);
}
starpu_cuda_set_device(devid);
}
#endif
switch (_STARPU_MEMORY_NODE_TUPLE(src_kind,dst_kind))
{
case _STARPU_MEMORY_NODE_TUPLE(STARPU_CPU_RAM,STARPU_CPU_RAM):
/* STARPU_CPU_RAM -> STARPU_CPU_RAM */
if (copy_methods->ram_to_ram)
copy_methods->ram_to_ram(src_interface, src_node, dst_interface, dst_node);
else
copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, req ? &req->async_channel : NULL);
break;
#ifdef STARPU_USE_CUDA
case _STARPU_MEMORY_NODE_TUPLE(STARPU_CUDA_RAM,STARPU_CPU_RAM):
/* only the proper CUBLAS thread can initiate this directly ! */
#if !defined(HAVE_CUDA_MEMCPY_PEER)
STARPU_ASSERT(_starpu_memory_node_get_local_key() == src_node);
#endif
if (!req || starpu_asynchronous_copy_disabled() || starpu_asynchronous_cuda_copy_disabled() ||
!(copy_methods->cuda_to_ram_async || copy_methods->any_to_any))
{
/* this is not associated to a request so it's synchronous */
STARPU_ASSERT(copy_methods->cuda_to_ram || copy_methods->any_to_any);
if (copy_methods->cuda_to_ram)
copy_methods->cuda_to_ram(src_interface, src_node, dst_interface, dst_node);
else
copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, NULL);
}
else
{
req->async_channel.type = STARPU_CUDA_RAM;
cures = cudaEventCreateWithFlags(&req->async_channel.event.cuda_event, cudaEventDisableTiming);
if (STARPU_UNLIKELY(cures != cudaSuccess)) STARPU_CUDA_REPORT_ERROR(cures);
stream = starpu_cuda_get_local_out_transfer_stream();
if (copy_methods->cuda_to_ram_async)
ret = copy_methods->cuda_to_ram_async(src_interface, src_node, dst_interface, dst_node, stream);
else
{
STARPU_ASSERT(copy_methods->any_to_any);
ret = copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, &req->async_channel);
}
cures = cudaEventRecord(req->async_channel.event.cuda_event, stream);
if (STARPU_UNLIKELY(cures != cudaSuccess)) STARPU_CUDA_REPORT_ERROR(cures);
}
break;
case _STARPU_MEMORY_NODE_TUPLE(STARPU_CPU_RAM,STARPU_CUDA_RAM):
/* STARPU_CPU_RAM -> CUBLAS_RAM */
/* only the proper CUBLAS thread can initiate this ! */
#if !defined(HAVE_CUDA_MEMCPY_PEER)
STARPU_ASSERT(_starpu_memory_node_get_local_key() == dst_node);
#endif
if (!req || starpu_asynchronous_copy_disabled() || starpu_asynchronous_cuda_copy_disabled() ||
!(copy_methods->ram_to_cuda_async || copy_methods->any_to_any))
{
/* this is not associated to a request so it's synchronous */
STARPU_ASSERT(copy_methods->ram_to_cuda || copy_methods->any_to_any);
if (copy_methods->ram_to_cuda)
copy_methods->ram_to_cuda(src_interface, src_node, dst_interface, dst_node);
else
copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, NULL);
}
else
{
req->async_channel.type = STARPU_CUDA_RAM;
cures = cudaEventCreateWithFlags(&req->async_channel.event.cuda_event, cudaEventDisableTiming);
if (STARPU_UNLIKELY(cures != cudaSuccess))
STARPU_CUDA_REPORT_ERROR(cures);
stream = starpu_cuda_get_local_in_transfer_stream();
if (copy_methods->ram_to_cuda_async)
ret = copy_methods->ram_to_cuda_async(src_interface, src_node, dst_interface, dst_node, stream);
else
{
STARPU_ASSERT(copy_methods->any_to_any);
ret = copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, &req->async_channel);
}
cures = cudaEventRecord(req->async_channel.event.cuda_event, stream);
if (STARPU_UNLIKELY(cures != cudaSuccess))
STARPU_CUDA_REPORT_ERROR(cures);
}
break;
case _STARPU_MEMORY_NODE_TUPLE(STARPU_CUDA_RAM,STARPU_CUDA_RAM):
/* CUDA - CUDA transfer */
if (!req || starpu_asynchronous_copy_disabled() || starpu_asynchronous_cuda_copy_disabled() ||
!(copy_methods->cuda_to_cuda_async || copy_methods->any_to_any))
{
STARPU_ASSERT(copy_methods->cuda_to_cuda || copy_methods->any_to_any);
/* this is not associated to a request so it's synchronous */
if (copy_methods->cuda_to_cuda)
copy_methods->cuda_to_cuda(src_interface, src_node, dst_interface, dst_node);
else
copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, NULL);
}
else
{
req->async_channel.type = STARPU_CUDA_RAM;
cures = cudaEventCreateWithFlags(&req->async_channel.event.cuda_event, cudaEventDisableTiming);
if (STARPU_UNLIKELY(cures != cudaSuccess)) STARPU_CUDA_REPORT_ERROR(cures);
stream = starpu_cuda_get_peer_transfer_stream(src_node, dst_node);
if (copy_methods->cuda_to_cuda_async)
ret = copy_methods->cuda_to_cuda_async(src_interface, src_node, dst_interface, dst_node, stream);
else
{
STARPU_ASSERT(copy_methods->any_to_any);
ret = copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, &req->async_channel);
}
cures = cudaEventRecord(req->async_channel.event.cuda_event, stream);
if (STARPU_UNLIKELY(cures != cudaSuccess)) STARPU_CUDA_REPORT_ERROR(cures);
}
break;
#endif
#ifdef STARPU_USE_OPENCL
case _STARPU_MEMORY_NODE_TUPLE(STARPU_OPENCL_RAM,STARPU_CPU_RAM):
/* OpenCL -> RAM */
STARPU_ASSERT(_starpu_memory_node_get_local_key() == src_node);
if (!req || starpu_asynchronous_copy_disabled() || starpu_asynchronous_opencl_copy_disabled() ||
!(copy_methods->opencl_to_ram_async || copy_methods->any_to_any))
{
STARPU_ASSERT(copy_methods->opencl_to_ram || copy_methods->any_to_any);
/* this is not associated to a request so it's synchronous */
if (copy_methods->opencl_to_ram)
copy_methods->opencl_to_ram(src_interface, src_node, dst_interface, dst_node);
else
copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, NULL);
}
else
{
req->async_channel.type = STARPU_OPENCL_RAM;
if (copy_methods->opencl_to_ram_async)
ret = copy_methods->opencl_to_ram_async(src_interface, src_node, dst_interface, dst_node, &(req->async_channel.event.opencl_event));
else
{
STARPU_ASSERT(copy_methods->any_to_any);
ret = copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, &req->async_channel);
}
}
break;
case _STARPU_MEMORY_NODE_TUPLE(STARPU_CPU_RAM,STARPU_OPENCL_RAM):
/* STARPU_CPU_RAM -> STARPU_OPENCL_RAM */
STARPU_ASSERT(_starpu_memory_node_get_local_key() == dst_node);
if (!req || starpu_asynchronous_copy_disabled() || starpu_asynchronous_opencl_copy_disabled() ||
!(copy_methods->ram_to_opencl_async || copy_methods->any_to_any))
{
STARPU_ASSERT(copy_methods->ram_to_opencl || copy_methods->any_to_any);
/* this is not associated to a request so it's synchronous */
if (copy_methods->ram_to_opencl)
copy_methods->ram_to_opencl(src_interface, src_node, dst_interface, dst_node);
else
copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, NULL);
}
else
{
req->async_channel.type = STARPU_OPENCL_RAM;
if (copy_methods->ram_to_opencl_async)
ret = copy_methods->ram_to_opencl_async(src_interface, src_node, dst_interface, dst_node, &(req->async_channel.event.opencl_event));
else
{
STARPU_ASSERT(copy_methods->any_to_any);
ret = copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, &req->async_channel);
}
}
break;
case _STARPU_MEMORY_NODE_TUPLE(STARPU_OPENCL_RAM,STARPU_OPENCL_RAM):
/* STARPU_OPENCL_RAM -> STARPU_OPENCL_RAM */
STARPU_ASSERT(_starpu_memory_node_get_local_key() == dst_node || _starpu_memory_node_get_local_key() == src_node);
if (!req || starpu_asynchronous_copy_disabled() || starpu_asynchronous_opencl_copy_disabled() ||
!(copy_methods->opencl_to_opencl_async || copy_methods->any_to_any))
{
STARPU_ASSERT(copy_methods->opencl_to_opencl || copy_methods->any_to_any);
/* this is not associated to a request so it's synchronous */
if (copy_methods->opencl_to_opencl)
copy_methods->opencl_to_opencl(src_interface, src_node, dst_interface, dst_node);
else
copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, NULL);
}
else
{
req->async_channel.type = STARPU_OPENCL_RAM;
if (copy_methods->opencl_to_opencl_async)
ret = copy_methods->opencl_to_opencl_async(src_interface, src_node, dst_interface, dst_node, &(req->async_channel.event.opencl_event));
else
{
STARPU_ASSERT(copy_methods->any_to_any);
ret = copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, &req->async_channel);
}
}
break;
#endif
#ifdef STARPU_USE_MIC
case _STARPU_MEMORY_NODE_TUPLE(STARPU_CPU_RAM,STARPU_MIC_RAM):
/* RAM -> MIC */
if (!req || starpu_asynchronous_copy_disabled() || starpu_asynchronous_mic_copy_disabled() ||
!(copy_methods->ram_to_mic_async || copy_methods->any_to_any))
{
/* this is not associated to a request so it's synchronous */
STARPU_ASSERT(copy_methods->ram_to_mic || copy_methods->any_to_any);
if (copy_methods->ram_to_mic)
copy_methods->ram_to_mic(src_interface, src_node, dst_interface, dst_node);
else
copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, NULL);
}
else
{
req->async_channel.type = STARPU_MIC_RAM;
if (copy_methods->ram_to_mic_async)
ret = copy_methods->ram_to_mic_async(src_interface, src_node, dst_interface, dst_node);
else
{
STARPU_ASSERT(copy_methods->any_to_any);
ret = copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, &req->async_channel);
}
_starpu_mic_init_event(&(req->async_channel.event.mic_event), dst_node);
}
break;
case _STARPU_MEMORY_NODE_TUPLE(STARPU_MIC_RAM,STARPU_CPU_RAM):
/* MIC -> RAM */
if (!req || starpu_asynchronous_copy_disabled() || starpu_asynchronous_mic_copy_disabled() ||
!(copy_methods->mic_to_ram_async || copy_methods->any_to_any))
{
/* this is not associated to a request so it's synchronous */
STARPU_ASSERT(copy_methods->mic_to_ram || copy_methods->any_to_any);
if (copy_methods->mic_to_ram)
copy_methods->mic_to_ram(src_interface, src_node, dst_interface, dst_node);
else
copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, NULL);
}
else
{
req->async_channel.type = STARPU_MIC_RAM;
if (copy_methods->mic_to_ram_async)
ret = copy_methods->mic_to_ram_async(src_interface, src_node, dst_interface, dst_node);
else
{
STARPU_ASSERT(copy_methods->any_to_any);
ret = copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, &req->async_channel);
}
_starpu_mic_init_event(&(req->async_channel.event.mic_event), src_node);
}
break;
#endif
#ifdef STARPU_USE_SCC
/* SCC RAM associated to the master process is considered as
* the main memory node. */
case _STARPU_MEMORY_NODE_TUPLE(STARPU_CPU_RAM,STARPU_SCC_RAM):
/* master private SCC RAM -> slave private SCC RAM */
if (copy_methods->scc_src_to_sink)
copy_methods->scc_src_to_sink(src_interface, src_node, dst_interface, dst_node);
else
copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, NULL);
break;
case _STARPU_MEMORY_NODE_TUPLE(STARPU_SCC_RAM,STARPU_CPU_RAM):
/* slave private SCC RAM -> master private SCC RAM */
if (copy_methods->scc_sink_to_src)
copy_methods->scc_sink_to_src(src_interface, src_node, dst_interface, dst_node);
else
copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, NULL);
break;
case _STARPU_MEMORY_NODE_TUPLE(STARPU_SCC_RAM,STARPU_SCC_RAM):
/* slave private SCC RAM -> slave private SCC RAM */
if (copy_methods->scc_sink_to_sink)
copy_methods->scc_sink_to_sink(src_interface, src_node, dst_interface, dst_node);
else
copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, NULL);
break;
#endif
case _STARPU_MEMORY_NODE_TUPLE(STARPU_CPU_RAM,STARPU_DISK_RAM):
if(copy_methods->any_to_any)
ret = copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, req && !starpu_asynchronous_copy_disabled() ? &req->async_channel : NULL);
else
{
void *obj = starpu_data_handle_to_pointer(handle, dst_node);
void * ptr = NULL;
starpu_ssize_t size = 0;
handle->ops->pack_data(handle, src_node, &ptr, &size);
ret = _starpu_disk_full_write(src_node, dst_node, obj, ptr, size, &req->async_channel);
if (ret == 0)
/* write is already finished, ptr was allocated in pack_data */
free(ptr);
/* For now, asynchronous is not supported */
STARPU_ASSERT(ret == 0);
}
break;
case _STARPU_MEMORY_NODE_TUPLE(STARPU_DISK_RAM,STARPU_CPU_RAM):
if(copy_methods->any_to_any)
ret = copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, req && !starpu_asynchronous_copy_disabled() ? &req->async_channel : NULL);
else
{
void *obj = starpu_data_handle_to_pointer(handle, src_node);
void * ptr = NULL;
size_t size = 0;
ret = _starpu_disk_full_read(src_node, dst_node, obj, &ptr, &size, &req->async_channel);
if (ret == 0)
{
/* read is already finished, we can already unpack */
handle->ops->unpack_data(handle, dst_node, ptr, size);
/* ptr is allocated in full_read */
free(ptr);
}
/* For now, asynchronous is not supported */
STARPU_ASSERT(ret == 0);
}
break;
case _STARPU_MEMORY_NODE_TUPLE(STARPU_DISK_RAM,STARPU_DISK_RAM):
ret = copy_methods->any_to_any(src_interface, src_node, dst_interface, dst_node, req ? &req->async_channel : NULL);
break;
default:
STARPU_ABORT();
break;
}
return ret;
#endif /* !SIMGRID */
}
extern "C" magma_int_t
magma_zbulge_applyQ_v2_m(
magma_int_t ngpu,
magma_side_t side,
magma_int_t NE, magma_int_t N,
magma_int_t NB, magma_int_t Vblksiz,
magmaDoubleComplex *E, magma_int_t lde,
magmaDoubleComplex *V, magma_int_t ldv,
magmaDoubleComplex *T, magma_int_t ldt,
magma_int_t *info)
{
//%===========================
//% local variables
//%===========================
magma_int_t Vm, Vn, mt, nt;
magma_int_t myrow, mycol, blkj, blki;
magma_int_t blkid,vpos,tpos;
magma_int_t firstrow, nbcolinvolvd;
magma_int_t versionL = 113;
magma_int_t versionR = 92;
magma_int_t Vchunksiz = 10;
*info=0;
/* Quick return */
if ( NE == 0 ) {
return *info;
}
if ( N == 0 ) {
return *info;
}
if ( NB == 0 ) {
return *info;
}
/* ==========================================
* some infos for developer
* Initialisation and checking nb of cores
* ==========================================*/
/* we have 2 algo for left (113 114) and 2 algo for right (91 92)
* which correspond to versionL versionR.
* They are very similar (detail explained in tech report and matlab code)
* however version 114 and 92 improve locality.
* while version 113 is used in case WNATZ=1 (construct Q2) which allow
* the construction to be done in an optimized way taking into
* consideration that the matrix is Identity so making less flops.
*
*/
// Initialize streaming and events
magma_device_sync();
magma_device_t orig_dev;
magma_getdevice( &orig_dev );
magma_queue_t orig_stream;
magmablasGetKernelStream( &orig_stream );
magma_int_t nbevents =2, nstream=2;
magma_queue_t streams[MagmaMaxGPUs][20];
magma_event_t myevent[MagmaMaxGPUs][20];
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
for( magma_int_t i = 0; i < nstream; ++i ) {
magma_queue_create( &streams[dev][i] );
}
for( magma_int_t i = 0; i < nbevents; ++i ) {
cudaEventCreateWithFlags(&myevent[dev][i],cudaEventDisableTiming);
}
}
// Azzam 21/11/2012
// NOTE THAT dwork was of size 2*NE*Vblksiz+...
// but I am thinking why not modifing it to NE*Vblksiz+...
// BUT NO because the 2* is used because of making 2 streams working and so
// they might be using dwork in parallel
magmaDoubleComplex *dE[MagmaMaxGPUs];
magmaDoubleComplex *dwork[MagmaMaxGPUs], *dwork0[MagmaMaxGPUs], *dwork1[MagmaMaxGPUs];
//magmaDoubleComplex *dwvt[MagmaMaxGPUs];
magmaDoubleComplex *dwvt0[MagmaMaxGPUs], *dwvt1[MagmaMaxGPUs];
magmaDoubleComplex *dT0[MagmaMaxGPUs], *dV0[MagmaMaxGPUs], *dT1[MagmaMaxGPUs], *dV1[MagmaMaxGPUs];
magma_int_t dev;
magma_int_t ldde = N;
magma_int_t lddv = ldv;
magma_int_t lddt = ldt;
magma_int_t ne_loc = magma_ceildiv(NE, ngpu);
if (ne_loc < 256)
ne_loc=256;
magma_int_t dwVTsiz = lddv*Vblksiz; // lddv*lddv + lddv*NE; // lddv*Vblksiz;
magma_int_t dworksiz = ne_loc*Vblksiz; // lddv*Vblksiz; // NE*Vblksiz;
ngpu = min(ngpu, magma_ceildiv(NE,ne_loc)); // Don't use GPU that will not have data.
// copy dE to GPUs
for (dev=0; dev < ngpu; ++dev) {
magma_setdevice( dev );
if (MAGMA_SUCCESS != magma_zmalloc( &dE[dev], ldde * ne_loc)) {
printf ("!!!! magma_zbulge_applyQ magma_alloc failed for: dE\n" );
exit(-1);
}
if (MAGMA_SUCCESS != magma_zmalloc( &dwork[dev], 2*dworksiz + 2*dwVTsiz + 2*Vchunksiz* (Vblksiz* (lddv+lddt)) )) {
printf ("!!!! magma_zbulge_applyQ magma_alloc failed for: dwork\n" );
exit(-1);
}
dwork0[dev] = dwork[dev]; // size = dworksiz;
dwork1[dev] = dwork0[dev] + dworksiz; // size = dworksiz;
dwvt0[dev] = dwork[dev] + 2*dworksiz; // size = dwVTsiz;
dwvt1[dev] = dwvt0[dev] + dwVTsiz; // size = dwVTsiz;
dV0[dev] = dwork[dev] + 2*dworksiz + 2*dwVTsiz;
dT0[dev] = dV0[dev] + Vchunksiz*Vblksiz*lddv;
dV1[dev] = dT0[dev] + Vchunksiz*Vblksiz*lddt;
dT1[dev] = dV1[dev] + Vchunksiz*Vblksiz*lddv;
magma_int_t ie_loc = min(ne_loc, NE - ne_loc*dev);
magma_zsetmatrix_async( N, ie_loc, E+lde*ne_loc*dev, lde, dE(dev, 0, 0), ldde, streams[dev][1] );
}
// make overlapped copy
magma_int_t ncpy = 0;
magma_int_t copyed=0, copyst=0;
magma_int_t blkcnt,nothing, mysiz, flip, vld,tld, locpos;
findVTsiz(N, NB, Vblksiz, &blkcnt, ¬hing);
flip = 0;
/* SIDE LEFT meaning apply E = Q*E = (q_1*q_2*.....*q_n) * E ==> so traverse Vs in reverse order (forward) from q_n to q_1
* Also E is splitten by row meaning each apply consist in a block of row (horizontal block) */
/* SIDE RIGHT meaning apply E = E*Q = E * (q_1*q_2*.....*q_n) ==> so tarverse Vs in normal order (forward) from q_1 to q_n
* Also E is splitten by col meaning each apply consist in a block of col (vertical block) */
#ifdef ENABLE_DEBUG
printf(" APPLY Q_v22_m GPU with NGPU %d N %d, NE %d, NB %d, Vblksiz %d, versionL %d versionR %d SIDE %c \n",
ngpu, N, NE, NB, Vblksiz, versionL, versionR, side);
#endif
/*
* MagmamaLeft
*/
if (side == MagmaLeft) {
/*
* Version 113:
* loop over the block_col (nt) and for each find the
* number of tiles (mt) in this block_col. then loop over mt, find
* the size of the V's(Vm,Vn) and apply it to the corresponding
* portion of E.
*/
if ( versionL == 113 ) {
nt = magma_ceildiv((N-1),Vblksiz);
for (blkj=nt-1; blkj >= 0; blkj--) {
/* the index of the first row on the top of block (blkj) */
firstrow = blkj * Vblksiz + 1;
/*find the number of tile for this block */
if ( blkj == nt-1 )
mt = magma_ceildiv( N - firstrow, NB);
else
mt = magma_ceildiv( N - (firstrow+1), NB);
/*loop over the tiles find the size of the Vs and apply it */
for (blki=mt; blki > 0; blki--) {
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + (mt-blki)*NB;
mycol = blkj*Vblksiz;
Vm = min( NB+Vblksiz-1, N-myrow);
if ( ( blkj == nt-1 ) && ( blki == mt ) ) {
Vn = min (Vblksiz, Vm);
} else {
Vn = min (Vblksiz, Vm-1);
}
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
//printf("voici blkj %d blki %d Vm %d Vn %d mycol %d vpos %d \n",blkj,blki,Vm, Vn,mycol,vpos);
magma_bulge_findpos113(N, NB, Vblksiz, mycol, myrow, &blkid);
// COPY Vchunksiz Vs and Vchunksiz Ts to GPU and store it in dV0/dV1 and dT0/dT1
if (ncpy == 0) {
// flip = 1 for this.
copyst = 0; // meaning that copy will start copying from blkid =copyst
copyed = min(copyst+Vchunksiz, blkcnt); // meaning that copy will end copying at blkid =copyed-1==> next copy had to start at copyed
mysiz = copyed-copyst; // the size of the chunk to be copied
if (mysiz > 0) {
ncpy = 1;
flip = 1;
vpos = copyst*Vblksiz*ldv;
tpos = copyst*Vblksiz*ldt;
vld = mysiz * ldv;
tld = mysiz * ldt;
for( dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 1 ] );
magma_zsetmatrix_async(vld, Vblksiz, V(vpos), vld, dV1[dev], vld, streams[dev][1]);
magma_zsetmatrix_async(tld, Vblksiz, T(tpos), tld, dT1[dev], tld, streams[dev][1]);
}
//printf("doing the first copy of mysiz %2d copyst %2d copyed %2d vpos %8d tpos %8d into dV1 dT1\n",mysiz,copyst,copyed,vpos,tpos);
}
}
if (blkid == copyst) {
flip = ncpy % 2;
copyst = copyed; // meaning that copy will start copying from blkid =copyst
copyed = min(copyst+Vchunksiz, blkcnt); // meaning that copy will end copying at blkid =copyed-1==> next copy had to start at copyed
mysiz = copyed-copyst; // the size of the chunk to be copied
//printf(" get to copy blkid %d blkid+(2*Vchunksiz) %d copyst %d copyed %d\n",blkid,blkid+(Vchunksiz),copyst,copyed);
if (mysiz > 0) {
ncpy = ncpy + 1;
vpos = copyst*Vblksiz*ldv;
tpos = copyst*Vblksiz*ldt;
vld = mysiz * ldv;
tld = mysiz * ldt;
if (flip == 0) { // now I am working on dV0 so copy the next and put it on dV1
//printf("doing overlapping copy of mysiz %2d copyst %2d copyed %2d vpos %8d tpos %8d into dV1 dT1\n",mysiz,copyst,copyed,vpos,tpos);
for( dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 1 ] );
magma_zsetmatrix_async(vld, Vblksiz, V(vpos), vld, dV1[dev], vld, streams[dev][1]);
magma_zsetmatrix_async(tld, Vblksiz, T(tpos), tld, dT1[dev], tld, streams[dev][1]);
}
} else { // now I am working on dV1 so copy the next and put it on dV0
//printf("doing overlapping copy of mysiz %2d copyst %2d copyed %2d vpos %8d tpos %8d into dV0 dT0\n",mysiz,copyst,copyed,vpos,tpos);
for( dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream( streams[ dev ][ 0 ] );
magma_zsetmatrix_async(vld, Vblksiz, V(vpos), vld, dV0[dev], vld, streams[dev][0]);
magma_zsetmatrix_async(tld, Vblksiz, T(tpos), tld, dT0[dev], tld, streams[dev][0]);
}
}
}
}
if ((Vm > 0) && (Vn > 0)) {
locpos = blkid%Vchunksiz;
magma_int_t lcvpos = locpos*Vblksiz*lddv;
magma_int_t lctpos = locpos*Vblksiz*lddt;
//printf("voici blkj %d blki %d Vm %d Vn %d mycol %d locvpos %5d loctpos %5d blkid %2d using data in dV%1d dT%1d \n",blkj,blki,Vm, Vn,mycol,lcvpos,lctpos, blkid,flip,flip);
if (flip == 0) {
for( dev = 0; dev < ngpu; ++dev ) {
magma_int_t ie_loc = min(ne_loc, NE - ne_loc*dev);
magma_int_t nr_bl = magma_ceildiv(ie_loc,10000); //nr of blocks
magma_int_t sz_bl = magma_ceildiv(ie_loc,nr_bl*64)*64; //maximum size of blocks (to have blocks of around the same size and multiple of 64)
magma_int_t ib; //size of current block
magma_setdevice( dev );
magmablasSetKernelStream(streams[dev][0]);
magma_queue_wait_event( streams[dev][0], myevent[dev][1] );
for (magma_int_t i=0; i < ie_loc; i += sz_bl) {
ib = min(sz_bl, ie_loc-i);
//magma_zlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV0[dev]+lcvpos, lddv, dT0[dev]+lctpos, lddt, dE(dev,myrow,i), ldde, dwork0[dev], ib);
magma_zlarfb_gpu_gemm( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV0[dev]+lcvpos, lddv, dT0[dev]+lctpos, lddt, dE(dev,myrow,i), ldde, dwork0[dev], ib, dwvt0[dev], Vm);
}
magma_event_record( myevent[dev][0], streams[dev][0] );
}
} else {
for( dev = 0; dev < ngpu; ++dev ) {
magma_int_t ie_loc = min(ne_loc, NE - ne_loc*dev);
magma_int_t nr_bl = magma_ceildiv(ie_loc,10000); //nr of blocks
magma_int_t sz_bl = magma_ceildiv(ie_loc,nr_bl*64)*64; //maximum size of blocks (to have blocks of around the same size and multiple of 64)
magma_int_t ib; //size of current block
magma_setdevice( dev );
magmablasSetKernelStream(streams[dev][1]);
magma_queue_wait_event( streams[dev][1], myevent[dev][0] );
for (magma_int_t i=0; i < ie_loc; i += sz_bl) {
ib = min(sz_bl, ie_loc-i);
//magma_zlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV1[dev]+lcvpos, lddv, dT1[dev]+lctpos, lddt, dE(dev,myrow,i), ldde, dwork1[dev], ib);
magma_zlarfb_gpu_gemm( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV1[dev]+lcvpos, lddv, dT1[dev]+lctpos, lddt, dE(dev,myrow,i), ldde, dwork1[dev], ib, dwvt1[dev], Vm);
}
magma_event_record( myevent[dev][1], streams[dev][1] );
}
}
} // end for (Vm &Vn) > 0
} // end for blki
} // end for blkj
} // end if version=113
/*
* Version 114:
* loop over the block_row (mt) and for each find diagonally the
* number of tiles (nt) in this block_row. then loop over nt, find
* the size of the V's(Vm,Vn) and apply it to the corresponding
* portion of E.
*/
else {
printf("versionL 114 not implemented in zbulge_applyQ_v2_m\n");
exit(-1);
mt = magma_ceildiv((N-1),NB);
for (blki = mt; blki > 0; blki--) {
/* nbcolinvolvd = number of column corresponding to this block_row (blki) */
nbcolinvolvd = min(N-1, blki*NB);
/*find the number of tile for this block (diagonal row of tiles) */
nt = magma_ceildiv(nbcolinvolvd,Vblksiz);
/*loop over the tiles find the size of the Vs and apply it */
for (blkj = nt-1; blkj >= 0; blkj--) {
/* the index of the first row of the first col meaning
* the block on the top left (blki) */
firstrow = (mt-blki)*NB+1;
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + blkj*Vblksiz;
mycol = blkj*Vblksiz;
Vm = min( NB+Vblksiz-1, N-myrow);
if ( ( blkj == nt-1 ) && ( blki == mt ) ) {
Vn = min (Vblksiz, Vm);
} else {
Vn = min (Vblksiz, Vm-1);
}
if ((Vm > 0) && (Vn > 0)) {
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
/*
magma_bulge_findVTpos(N, NB, Vblksiz, mycol, myrow, ldv, ldt, &vpos, &tpos);
magma_zsetmatrix_async(Vm, Vn, V(vpos), ldv, dV0, lddv, NULL);
magma_zsetmatrix_async(Vn, Vn, T(tpos), ldt, dT0, lddt, NULL);
//printf("voici blki %d rownbm %d mycol %d coled %d blkid %d vpos %d tpos %d\n", blki, rownbm, mycol, coled, blkid, vpos, tpos);
for (magma_int_t i=0; i < NE; i += sz_bl) {
ib = min(sz_bl, NE-i);
magma_zlarfb_gpu( MagmaLeft, MagmaNoTrans, MagmaForward, MagmaColumnwise, Vm, ib, Vn, dV0, lddv, dT0, lddt, dE(myrow,i), ldde, dwork, NE);
}
*/
} // end for (Vm &Vn) > 0
} // end for blkj
} // end for blki
} // end version 114
} // end LEFT
/*
* MagmaRight
*/
else {
printf("Side 'R' not implemented in zbulge_applyQ_v2_m\n");
exit(-1);
/*
* Version 91:
*/
if ( versionR == 91 ) {
nt = magma_ceildiv((N-1),Vblksiz);
for (blkj=0; blkj < nt; blkj++) {
/* the index of the first myrow on the top of block (blkj) */
firstrow = blkj * Vblksiz + 1;
/*find the number of tile for this block */
if ( blkj == nt-1 )
mt = magma_ceildiv( N - firstrow, NB);
else
mt = magma_ceildiv( N - (firstrow+1), NB);
/*loop over the tiles find the size of the Vs and apply it */
for (blki=1; blki <= mt; blki++) {
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + (mt-blki)*NB;
Vm = min( NB+Vblksiz-1, N-myrow);
if ( (blkj == nt-1) && (blki == mt) ) {
Vn = min (Vblksiz, Vm);
} else {
Vn = min (Vblksiz, Vm-1);
}
mycol = blkj*Vblksiz;
if ((Vm > 0) && (Vn > 0)) {
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
/*
magma_bulge_findVTpos(N, NB, Vblksiz, mycol, myrow, ldv, ldt, &vpos, &tpos);
magma_zsetmatrix_async(Vm, Vn, V(vpos), ldv, dV0, lddv, NULL);
magma_zsetmatrix_async(Vn, Vn, T(tpos), ldt, dT0, lddt, NULL);
magma_zlarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, NE, Vm, Vn, dV0, lddv, dT0, lddt, dE(0, myrow), ldde, dwork, NE);
*/
} // end for (Vm &Vn) > 0
} // end for blki
} // end fo blkj
} // end of version 91
/*
* Version 92:
*/
else {
mt = magma_ceildiv((N-1),NB);
for (blki = 1; blki <= mt; blki++) {
/* nbcolinvolvd = number of column corresponding to this block_row (blki) */
nbcolinvolvd = min(N-1, blki*NB);
/*find the number of tile for this block (diagonal row of tiles) */
nt = magma_ceildiv(nbcolinvolvd,Vblksiz);
/*loop over the tiles find the size of the Vs and apply it */
for (blkj = 0; blkj < nt; blkj++) {
/* the index of the first row of the first col meaning
* the block on the top left (blki) */
firstrow = (mt-blki)*NB+1;
/*calculate the size of each losange of Vs= (Vm,Vn)*/
myrow = firstrow + blkj*Vblksiz;
mycol = blkj*Vblksiz;
Vm = min( NB+Vblksiz-1, N-myrow);
if ( ( blkj == nt-1 ) && ( blki == mt ) ) {
Vn = min (Vblksiz, Vm);
} else {
Vn = min (Vblksiz, Vm-1);
}
if ((Vm > 0) && (Vn > 0)) {
/*calculate the pointer to the Vs and the Ts.
* Note that Vs and Ts have special storage done
* by the bulgechasing function*/
/*
magma_bulge_findVTpos(N, NB, Vblksiz, mycol, myrow, ldv, ldt, &vpos, &tpos);
magma_zsetmatrix_async(Vm, Vn, V(vpos), ldv, dV0, lddv, NULL);
magma_zsetmatrix_async(Vn, Vn, T(tpos), ldt, dT0, lddt, NULL);
magma_zlarfb_gpu( MagmaRight, MagmaNoTrans, MagmaForward, MagmaColumnwise, NE, Vm, Vn, dV0, lddv, dT0, lddt, dE(0, myrow), ldde, dwork, NE);
*/
} // end for (Vm &Vn) > 0
} //end for blkj
} // end for blki
} //end of version 92
} // end RIGHT
// copy back the dE form each GPU
for( dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream(streams[dev][0]);
magma_queue_wait_event( streams[dev][0], myevent[dev][1] );
magma_queue_wait_event( streams[dev][0], myevent[dev][0] );
magma_int_t ie_loc = min(ne_loc, NE - ne_loc*dev);
magma_zgetmatrix_async( N, ie_loc, dE(dev, 0, 0), ldde, E+lde*ne_loc*dev, lde, streams[dev][0] );
magma_event_record( myevent[dev][0], streams[dev][0] );
}
for( magma_int_t dev = 0; dev < ngpu; ++dev ) {
magma_setdevice( dev );
magmablasSetKernelStream(streams[dev][0]);
magma_queue_wait_event( streams[dev][0], myevent[dev][0] );
magma_device_sync(); // no need for synchronize
magma_free(dwork[dev]);
magma_free(dE[dev]);
for( magma_int_t i = 0; i < nbevents; ++i ) {
magma_event_destroy( myevent[dev][i] );
}
for( magma_int_t i = 0; i < nstream; ++i ) {
magma_queue_destroy( streams[dev][i] );
}
}
magma_setdevice( orig_dev );
magmablasSetKernelStream( orig_stream );
return *info;
}
/* ////////////////////////////////////////////////////////////////////////////
-- Testing magma_dsymm_mgpu
*/
int main( int argc, char** argv)
{
TESTING_INIT();
double c_neg_one = MAGMA_D_NEG_ONE;
double alpha = MAGMA_D_MAKE( 3.456, 5.678 );
double beta = MAGMA_D_MAKE( 1.234, 2.456 );
real_Double_t gflops, gpu_perf=0., cpu_perf=0., gpu_time=0., cpu_time=0.;
real_Double_t gpu_perf2=0., gpu_time2=0.;
double Anorm, error, work[1];
double *hA, *hB, *hC, *hR;
magmaDouble_ptr dA[MagmaMaxGPUs], dB[MagmaMaxGPUs], dC[MagmaMaxGPUs], dwork[MagmaMaxGPUs];
magmaDouble_ptr dA2;
magma_int_t i, j, dev, M, N, size, lda, ldb, ldc, ldda, lddb, lddc, msize, nb;
magma_int_t ione = 1;
magma_int_t iseed[4] = {0,0,0,1};
magma_int_t status = 0;
magma_opts opts;
opts.parse_opts( argc, argv );
opts.ngpu = abs( opts.ngpu ); // always uses multi-GPU code
double tol = opts.tolerance * lapackf77_dlamch("E");
// default values
nb = (opts.nb > 0 ? opts.nb : 64);
magma_int_t gnode[MagmaMaxGPUs][MagmaMaxGPUs+2];
magma_int_t ncmplx = 0;
magma_buildconnection_mgpu( gnode, &ncmplx, opts.ngpu );
printf("%% Initializing communication pattern... GPU-ncmplx %d\n", (int) ncmplx);
for (i=0; i < ncmplx; ++i) {
magma_int_t myngpu = gnode[i][MagmaMaxGPUs];
printf("%% cmplx %d has %d GPUs:", i, myngpu);
for (j=0; j < myngpu; ++j) {
printf(" %d", (int) gnode[i][j]);
if (j < myngpu-1) {
printf(",");
}
}
printf("\n");
}
// number of queues per GPU. Requires ngpu.
magma_int_t nqueue = opts.ngpu;
// number of events per GPU. Require ngpu*ngpu.
magma_int_t nevents = opts.ngpu*opts.ngpu;
magma_queue_t queues[MagmaMaxGPUs][20], queues0[MagmaMaxGPUs];
magma_event_t events[MagmaMaxGPUs][MagmaMaxGPUs*MagmaMaxGPUs + 10];
for( dev = 0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
for( i = 0; i < nqueue; ++i ) {
magma_queue_create( dev, &queues[dev][i] );
}
queues0[dev] = queues[dev][0];
for( i = 0; i < nevents; ++i ) {
cudaEventCreateWithFlags( &events[dev][i], cudaEventDisableTiming );
}
}
printf("%% nb %d, ngpu %d, version %d\n", (int) nb, (int) opts.ngpu, (int) opts.version );
printf("%% M N nb offset CPU Gflop/s (sec) GPU Gflop/s (sec) CUBLAS hemm (sec) ||R|| / ||A||*||B||\n");
printf("%%========================================================================================================\n");
for( int itest = 0; itest < opts.ntest; ++itest ) {
M = opts.msize[itest];
N = opts.nsize[itest];
for( int offset = 0; offset < N; offset += min(N,nb) ) {
for( int iter = 0; iter < opts.niter; ++iter ) {
msize = M - offset;
lda = M; // TODO depends on side
ldb = M;
ldc = M;
ldda = magma_roundup( lda, opts.align ); // multiple of 32 by default
lddb = magma_roundup( ldb, opts.align ); // multiple of 32 by default
lddc = magma_roundup( ldc, opts.align ); // multiple of 32 by default
gflops = FLOPS_DSYMM( MagmaLeft, (double)msize, (double)N ) / 1e9;
magma_int_t dworksiz = lddc*N + (M*N)*opts.ngpu;
TESTING_MALLOC_CPU( hA, double, lda*M );
TESTING_MALLOC_CPU( hB, double, ldb*N );
TESTING_MALLOC_CPU( hC, double, ldc*N );
TESTING_MALLOC_PIN( hR, double, ldc*N );
for( dev = 0; dev < opts.ngpu; ++dev ) {
magma_int_t mlocal = ((M / nb) / opts.ngpu + 1) * nb;
magma_setdevice( dev );
TESTING_MALLOC_DEV( dA[dev], double, ldda*mlocal );
TESTING_MALLOC_DEV( dB[dev], double, lddb*N );
TESTING_MALLOC_DEV( dC[dev], double, lddc*N );
TESTING_MALLOC_DEV( dwork[dev], double, dworksiz );
}
if ( opts.check ) {
magma_setdevice( 0 );
TESTING_MALLOC_DEV( dA2, double, ldda*M );
}
size = lda*M;
lapackf77_dlarnv( &ione, iseed, &size, hA );
magma_dmake_symmetric( M, hA, lda );
size = ldb*N;
lapackf77_dlarnv( &ione, iseed, &size, hB );
size = ldc*N;
lapackf77_dlarnv( &ione, iseed, &size, hC );
lapackf77_dlacpy( "Full", &M, &N, hC, &ldc, hR, &lda );
/* ====================================================================
Performs operation using MAGMA
=================================================================== */
magma_dsetmatrix_1D_col_bcyclic( M, M, hA, lda, dA, ldda, opts.ngpu, nb, queues0 );
for( dev = 0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
magma_dsetmatrix( M, N, hB, lda, dB[dev], ldda, opts.queue );
// since when offset != 0, the GPU that does beta*C may not be 0,
// send initial hC to all GPUs.
magma_dsetmatrix( M, N, hC, lda, dC[dev], ldda, opts.queue );
}
trace_init( 1, opts.ngpu, nqueue, (magma_queue_t*) queues );
gpu_time = magma_sync_wtime(0);
magmablas_dsymm_mgpu(
MagmaLeft, MagmaLower, msize, N,
alpha, dA, ldda, offset,
dB, ldda,
beta, dC, ldda, dwork, dworksiz,
opts.ngpu, nb, queues, nqueue, events, nevents, gnode, ncmplx);
gpu_time = magma_sync_wtime(0) - gpu_time;
gpu_perf = gflops / gpu_time;
#ifdef TRACING
char buf[80];
snprintf( buf, sizeof(buf), "dsymm-m%d-n%d-nb%d-ngpu%d-run%d.svg",
(int) M, (int) N, (int) nb, (int) opts.ngpu, (int) iter );
trace_finalize( buf, "trace.css" );
#endif
/* ====================================================================
Performs operation using CUBLAS
=================================================================== */
if ( opts.check && iter == 0 ) {
magma_setdevice( 0 );
magma_dsetmatrix( M, M, hA, lda, dA2, ldda, opts.queue );
magma_dsetmatrix( M, N, hB, lda, dB[0], ldda, opts.queue );
magma_dsetmatrix( M, N, hC, lda, dwork[0], ldda, opts.queue );
gpu_time2 = magma_sync_wtime(0);
magma_dsymm(
MagmaLeft, MagmaLower, msize, N,
alpha, dA2 + offset + offset*ldda, ldda,
dB[0], ldda,
beta, dwork[0], ldda, opts.queue );
gpu_time2 = magma_sync_wtime(0) - gpu_time2;
gpu_perf2 = gflops / gpu_time2;
}
/* =====================================================================
Performs operation using LAPACK
=================================================================== */
if ( opts.check ) {
// store ||A||*||B||
Anorm = lapackf77_dlange("fro", &msize, &msize, hA + offset + offset*lda, &lda, work );
Anorm *= lapackf77_dlange("fro", &msize, &N, hB, &lda, work );
//printf( "A =" ); magma_dprint( M, M, hA, lda );
//printf( "B =" ); magma_dprint( M, N, hB, lda );
//printf( "C =" ); magma_dprint( M, N, hC, lda );
cpu_time = magma_wtime();
blasf77_dsymm( "Left", "Lower", &msize, &N,
&alpha, hA + offset + offset*lda, &lda,
hB, &lda,
&beta, hC, &lda );
cpu_time = magma_wtime() - cpu_time;
cpu_perf = gflops / cpu_time;
for (dev=0; dev < opts.ngpu; ++dev) {
magma_setdevice( dev );
magma_dgetmatrix( M, N, dC[dev], ldda, hR, lda, opts.queue );
// compute relative error ||R||/||A||*||B||, where R := C_magma - C_lapack = R - C
size = ldc*N;
blasf77_daxpy( &size, &c_neg_one, hC, &ione, hR, &ione );
error = lapackf77_dlange("fro", &msize, &N, hR, &lda, work) / Anorm;
//printf( "R =" ); magma_dprint( M, N, hR, lda );
bool okay = (error < tol);
status += ! okay;
if (dev == 0) {
printf( "%5d %5d %5d %5d %7.1f (%7.4f) %7.1f (%7.4f) %7.1f (%7.4f) %8.2e %s\n",
(int) M, (int) N, (int) nb, (int) offset,
cpu_perf, cpu_time,
gpu_perf, gpu_time,
gpu_perf2, gpu_time2,
error, (okay ? "ok" : "failed") );
}
else {
printf( " dev %d %74s %8.2e %s\n", dev, "",
error, (okay ? "ok" : "failed") );
}
}
} else {
printf( "%5d %5d %5d %5d --- ( --- ) %7.1f (%7.4f) --- ( --- ) ---\n",
(int) M, (int) N, (int) nb, (int) offset,
gpu_perf, gpu_time );
}
TESTING_FREE_CPU( hA );
TESTING_FREE_CPU( hB );
TESTING_FREE_CPU( hC );
TESTING_FREE_PIN( hR );
for( dev = 0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
TESTING_FREE_DEV( dA[dev] );
TESTING_FREE_DEV( dB[dev] );
TESTING_FREE_DEV( dC[dev] );
TESTING_FREE_DEV( dwork[dev] );
}
if ( opts.check ) {
magma_setdevice( 0 );
TESTING_FREE_DEV( dA2 );
}
fflush( stdout );
}
if ( opts.niter > 1 ) {
printf( "\n" );
}
} // offset
printf( "\n" );
}
for( dev = 0; dev < opts.ngpu; ++dev ) {
magma_setdevice( dev );
for( i = 0; i < nqueue; ++i ) {
magma_queue_destroy( queues[dev][i] );
}
for( i = 0; i < nevents; ++i ) {
magma_event_destroy( events[dev][i] );
}
}
opts.cleanup();
TESTING_FINALIZE();
return status;
}
