/*
* Returns flops per cycle.
*/
double time_mult(long dim, enum mult_flag_enum flag){
int count = 10;
double *a = new double[dim*dim];
double *b = new double[dim*dim];
double *c = new double[dim*dim];
#pragma omp parallel for
for(long i=0; i < dim*dim; i++)
a[i] = b[i] = c[i] = 1;
TimeStamp clk;
StatVector stats(count);
if(flag == AUTO)
mkl_mic_enable();
for(int i=0; i < count; i++){
clk.tic();
switch(flag){
case HOST:
mmult(a, b, c, dim);
break;
case MIC:
#pragma offload target(mic) \
in(a:length(dim*dim) align(64) alloc_if(1) free_if(1)) \
in(b:length(dim*dim) align(64) alloc_if(1) free_if(1)) \
inout(c: length(dim*dim) align(64) alloc_if(1) free_if(1))
mmult(a, b, c, dim);
break;
case AUTO:
mmult(a, b, c, dim);
break;
}
double cycles = clk.toc();
stats.insert(cycles);
}
if(flag == AUTO)
mkl_mic_disable();
delete[] a;
delete[] b;
delete[] c;
return 2.0*dim*dim*dim/stats.median();
}
void cholesky_tiled(double *mat, int tile_size, int num_tiles, int mat_size,
int niter, int max_log_str, bool layRow, int verify)
{
//total number of tiles
int tot_tiles = num_tiles * num_tiles;
//memory allocation for matrix for tiled-Cholesky
double *A_my = (double *)malloc(mat_size * mat_size * sizeof(double));
//memory allocation for matrix for MKL cholesky (for comparison)
double *A_MKL = (double *)malloc(mat_size * mat_size * sizeof(double));
//memory allocation for tiled matrix
double **Asplit = new double* [tot_tiles];
int mem_size_tile = tile_size * tile_size * sizeof(double);
for (int i = 0; i < tot_tiles; ++i) {
//Buffer per tile, host allocation
Asplit[i] = (double *)_mm_malloc(mem_size_tile, 64);
//Buffer creation and allocation on the card
hStreams_app_create_buf((void *)Asplit[i], mem_size_tile);
}
double tbegin, tend;
int iter;
int info;
//Events are needed for various synchronizations to enforce
//data dependence between and among data-transfers/computes
HSTR_EVENT *eventcpyto = new HSTR_EVENT[tot_tiles];
HSTR_EVENT *eventcpyfr = new HSTR_EVENT[tot_tiles];
HSTR_EVENT *eventpotrf = new HSTR_EVENT[tot_tiles];
HSTR_EVENT *eventtrsm = new HSTR_EVENT[tot_tiles];
HSTR_EVENT *eventsyrk = new HSTR_EVENT[tot_tiles];
HSTR_EVENT *eventgemm = new HSTR_EVENT[tot_tiles];
//for timing tiled cholesky
double *totTimeMsec = new double [niter];
//for timing MKL cholesky
double *totTimeMsecMKL = new double [niter];
HSTR_RESULT res;
//these queues are used for queining up compute on the card and
//data transfers to/from the card.
//q_trsm for dtrsm, q_potrf for dportf, q_syrk_gemm for both dsyrk and dgemm.
//The queues are incremented by one for every compute queued and wrap
//around the max_log_str available. This ensures good load-balancing.
int q_trsm, q_potrf, q_syrk_gemm;
CBLAS_ORDER blasLay;
int lapackLay;
if (layRow) {
blasLay = CblasRowMajor;
lapackLay = LAPACK_ROW_MAJOR;
} else {
blasLay = CblasColMajor;
lapackLay = LAPACK_COL_MAJOR;
}
for (iter = 0; iter < niter; ++iter) {
//copying matrices into separate variables for tiled cholesky (A_my)
//and MKL cholesky (A_MKL)
//The output overwrites the matrices and hence the need to copy
//for each iteration
copy_mat(mat, A_my, mat_size);
copy_mat(mat, A_MKL, mat_size);
unsigned int m, n, k;
printf("\nIteration = %d\n", iter);
split_into_blocks(A_my, Asplit, num_tiles, tile_size, mat_size, layRow);
//beginning of timing
tbegin = dtimeGet();
//splitting time included in the timing
//This splits the input matrix into tiles (or blocks)
//split_into_blocks(A_my, Asplit, num_tiles, tile_size, mat_size, layRow);
q_potrf = 0;
for (k = 0; k < num_tiles; ++k) {
//POTRF
//dpotrf is executed on the host on the diagonal tile
//the results are then sent to the card
if (k > 0) {
hStreams_app_event_wait(1, &eventsyrk[k * num_tiles + k]);
if (loc_verbose > 0)
printf("Sending tile[%d][%d] to host in queue %d\n",
k, k, (int)(q_potrf % max_log_str)) ;
hStreams_app_xfer_memory(Asplit[k * num_tiles + k],
Asplit[k * num_tiles + k], mem_size_tile,
(int)(q_potrf % max_log_str), HSTR_SINK_TO_SRC,
&eventcpyfr[k * num_tiles + k]);
hStreams_app_event_wait(1, &eventcpyfr[k * num_tiles + k]);
}
if (loc_verbose > 0) {
printf("Executing potrf on host for tile[%d][%d]\n", k, k);
}
info = LAPACKE_dpotrf(lapackLay, 'L', tile_size,
Asplit[k * num_tiles + k], tile_size);
if (k < num_tiles - 1) {
if (loc_verbose > 0)
printf("Sending tile[%d][%d] to card in queue %d\n",
k, k, (int)(q_potrf % max_log_str));
hStreams_app_xfer_memory(Asplit[k * num_tiles + k],
Asplit[k * num_tiles + k], mem_size_tile,
(int)(q_potrf % max_log_str), HSTR_SRC_TO_SINK,
&eventcpyto[k * num_tiles + k]);
}
q_potrf++;
q_trsm = 0;
for (m = k + 1; m < num_tiles; ++m) {
if (k == 0) {
if (loc_verbose > 0)
printf("Sending tile[%d][%d] to card in queue %d\n",
m, k, (int)(q_trsm % max_log_str));
hStreams_app_xfer_memory(Asplit[m * num_tiles + k],
Asplit[m * num_tiles + k], mem_size_tile,
(int)(q_trsm % max_log_str), HSTR_SRC_TO_SINK,
&eventcpyto[m * num_tiles + k]);
}
//DTRSM
hStreams_app_event_wait(1, &eventcpyto[k * num_tiles + k]);
if (k > 0) {
hStreams_app_event_wait(1, &eventgemm[m * num_tiles + k]);
}
//dtrsm is executed on the card
if (loc_verbose > 0)
printf("Executing trsm for tile[%d][%d] on card in queue %d\n",
m, k, (int)(q_trsm % max_log_str));
res = hStreams_custom_dtrsm(blasLay, CblasRight, CblasLower,
CblasTrans, CblasNonUnit, tile_size, tile_size, 1.0,
Asplit[k * num_tiles + k], tile_size, Asplit[m * num_tiles + k],
tile_size, (int)(q_trsm % max_log_str),
&eventtrsm[m * num_tiles + k]);
if (loc_verbose > 0)
printf("Sending tile[%d][%d] back to host in queue %d\n",
m, k, (int)(q_trsm % max_log_str));
hStreams_app_xfer_memory(Asplit[m * num_tiles + k],
Asplit[m * num_tiles + k], mem_size_tile,
(int)(q_trsm % max_log_str), HSTR_SINK_TO_SRC,
&eventcpyfr[m * num_tiles + k]);
q_trsm++;
}
q_syrk_gemm = 0;
for (n = k + 1; n < num_tiles; ++n) {
if (k == 0) {
if (loc_verbose > 0)
printf("Sending tile[%d][%d] to card in queue %d\n",
n, n, (int)(q_syrk_gemm % max_log_str));
hStreams_app_xfer_memory(Asplit[n * num_tiles + n],
Asplit[n * num_tiles + n], mem_size_tile,
(int)(q_syrk_gemm % max_log_str), HSTR_SRC_TO_SINK,
&eventcpyto[n * num_tiles + n]);
}
//DSYRK
hStreams_app_event_wait(1, &eventtrsm[n * num_tiles + k]);
if (k > 0) {
hStreams_app_event_wait(1, &eventsyrk[n * num_tiles + n]);
}
//dsyrk is executed on the card
if (loc_verbose > 0)
printf("Executing syrk for tile[%d][%d] on card in queue %d\n",
n, n, (int)(q_syrk_gemm % max_log_str));
res = hStreams_custom_dsyrk(blasLay, CblasLower, CblasNoTrans,
tile_size, tile_size, -1.0, Asplit[n * num_tiles + k],
tile_size, 1.0, Asplit[n * num_tiles + n], tile_size,
(int)(q_syrk_gemm % max_log_str), &eventsyrk[n * num_tiles + n]);
q_syrk_gemm++;
for (m = n + 1; m < num_tiles; ++m) {
if (k == 0) {
if (loc_verbose > 0)
printf("Sending tile[%d][%d] to card in queue %d\n",
m, n, (int)(q_syrk_gemm % max_log_str));
hStreams_app_xfer_memory(Asplit[m * num_tiles + n],
Asplit[m * num_tiles + n], mem_size_tile,
(int)(q_syrk_gemm % max_log_str),
HSTR_SRC_TO_SINK,
&eventcpyto[m * num_tiles + n]);
}
//DGEMM
hStreams_app_event_wait(1, &eventtrsm[m * num_tiles + k]);
hStreams_app_event_wait(1, &eventtrsm[n * num_tiles + k]);
if (k > 0) {
hStreams_app_event_wait(1, &eventgemm[m * num_tiles + n]);
}
//dgemm is executed on the card
if (loc_verbose > 0)
printf("Executing gemm for tile[%d][%d] on card in queue %d\n",
m, n, (int)(q_syrk_gemm % max_log_str));
res = hStreams_app_dgemm(blasLay, CblasNoTrans, CblasTrans,
tile_size, tile_size, tile_size, -1.0, Asplit[m * num_tiles + k],
tile_size, Asplit[n * num_tiles + k], tile_size, 1.0,
Asplit[m * num_tiles + n], tile_size,
(int)(q_syrk_gemm % max_log_str), &eventgemm[m * num_tiles + n]);
q_syrk_gemm++;
}
}
}
//syncrhonizing all the streams
hStreams_app_thread_sync();
//end of timing
tend = dtimeGet();
totTimeMsec[iter] = 1e3 * (tend - tbegin);
printf("time for Tiled hstreams Cholesky for iteration %d = %.2f msec\n",
iter, totTimeMsec[iter]);
//assembling of tiles back into full matrix
assemble(Asplit, A_my, num_tiles, tile_size, mat_size, layRow);
//calling mkl cholesky for verification and timing comparison.
//Using auto-offload feature of MKL
#ifndef _WIN32
//FIXME: calling this function causes a crash on Windows
mkl_mic_enable();
#endif
tbegin = dtimeGet();
//calling MKL dpotrf on the full matrix
info = LAPACKE_dpotrf(lapackLay, 'L', mat_size, A_MKL, mat_size);
tend = dtimeGet();
totTimeMsecMKL[iter] = 1e3 * (tend - tbegin);
printf("time for MKL Cholesky (AO) for iteration %d = %.2f msec\n",
iter, totTimeMsecMKL[iter]);
if (info != 0) {
printf("error with dpotrf\n");
}
mkl_mic_disable();
if (verify == 1) {
bool result = verify_results(A_my, A_MKL, mat_size * mat_size);
if (result == true) {
printf("Tiled Cholesky successful\n");
} else {
printf("Tiled Chloesky failed\n");
}
}
}
double meanTimeMsec, stdDevMsec;
double meanTimeMsecMKL, stdDevMsecMKL;
mean_and_stdev(totTimeMsec, meanTimeMsec, stdDevMsec, niter);
mean_and_stdev(totTimeMsecMKL, meanTimeMsecMKL, stdDevMsecMKL, niter);
double gflops = pow(mat_size, 3.0) / 3.0 * 1e-9;
printf("\nMatrix size = %d\n", mat_size);
printf("Tiled hStreams Cholesky: for %d iterations (ignoring first),\n"
"mean Time = %.2f msec, stdDev Time = %.2f msec,\n"
"Mean Gflops (using mean Time) = %.2f\n",
niter - 1, meanTimeMsec, stdDevMsec, gflops / (meanTimeMsec * 1e-3));
printf("\nMKL AO Cholesky: for %d iterations (ignoring first),\n"
"mean Time = %.2f msec, stdDev Time = %.2f msec,\n"
"Mean Gflops (using meanTime) = %.2f\n\n",
niter - 1, meanTimeMsecMKL, stdDevMsecMKL, gflops / (meanTimeMsecMKL * 1e-3));
//Free
free(A_my);
free(A_MKL);
for (int i = 0; i < tot_tiles; ++i) {
_mm_free(Asplit[i]);
}
delete [] Asplit;
delete [] eventcpyto;
delete [] eventcpyfr;
delete [] eventpotrf;
delete [] eventtrsm;
delete [] eventsyrk;
delete [] eventgemm;
delete [] totTimeMsec;
delete [] totTimeMsecMKL;
}
