mirror of
https://github.com/cp2k/cp2k.git
synced 2026-07-28 14:15:19 -04:00
DBM/OFFLOAD: prepared for hybrid execution
- Split backend_upload_packs into two phases on GPU. - Initial phase is chosen to have not triggered upload. - Offload: offloadEventQuery to check if GPU is busy. - Removed dbm_multiply_gpu_download_results - Transfers are launched as early as possible. - Improved overlap between transfer/compute. - Idea: proceed on CPU if GPU is busy. Other changes: - Documented way to perform validation (ctx=NULL). - Permit DBM_ prefix for preprocessor symbols. - Removed C-shard pointer from GPU-context. - Per-shard event stored/provided by shard. - Migrated CPU-options into context. - Miniapp: allow to omit validation. - Use preincrement (atomic).
This commit is contained in:
Hans Pabst
parent
63a7c4bee4
commit
7face5689e
6 changed files with 184 additions and 181 deletions
|
|
@ -196,10 +196,15 @@ void benchmark_multiply(const int M, const int N, const int K, const int m,
|
|||
set_all_blocks(matrix_a);
|
||||
set_all_blocks(matrix_b);
|
||||
|
||||
dbm_distribution_t *const dist_shared = matrix_c->dist;
|
||||
dbm_create(&matrix_d, dist_shared, matrix_c->name, matrix_c->nrows,
|
||||
matrix_c->ncols, matrix_c->row_sizes, matrix_c->col_sizes);
|
||||
dbm_copy(matrix_d, matrix_c);
|
||||
const char *const verify_env = getenv("DBM_MULTIPLY_VERIFY");
|
||||
const int skip_verify = (NULL == verify_env ? 0 : (atoi(verify_env) + 1));
|
||||
|
||||
if (0 == skip_verify) {
|
||||
dbm_distribution_t *const dist_shared = matrix_c->dist;
|
||||
dbm_create(&matrix_d, dist_shared, matrix_c->name, matrix_c->nrows,
|
||||
matrix_c->ncols, matrix_c->row_sizes, matrix_c->col_sizes);
|
||||
dbm_copy(matrix_d, matrix_c);
|
||||
}
|
||||
|
||||
int64_t flop = 0;
|
||||
const double time_start_multiply = omp_get_wtime();
|
||||
|
|
@ -207,33 +212,33 @@ void benchmark_multiply(const int M, const int N, const int K, const int m,
|
|||
1e-8, &flop);
|
||||
const double time_end_multiply = omp_get_wtime();
|
||||
|
||||
// Calculate result on the host for validation
|
||||
dbm_multiply(false, false, 1.0, matrix_a, matrix_b, 1.0, matrix_d, false,
|
||||
1e-8, NULL);
|
||||
|
||||
if (dbm_mpi_comm_rank(comm) == 0) {
|
||||
printf("%5i x %5i x %5i with %3i x %3i x %3i blocks: ", M, N, K, m, n, k);
|
||||
}
|
||||
|
||||
// Validate result
|
||||
const double maxeps = 1E-5, epsilon = dbm_maxeps(matrix_d, matrix_c);
|
||||
if (maxeps >= epsilon) {
|
||||
dbm_mpi_sum_int64(&flop, 1, comm);
|
||||
if (dbm_mpi_comm_rank(comm) == 0) {
|
||||
const double duration = time_end_multiply - time_start_multiply;
|
||||
printf("%6.3f s => %6.1f GFLOP/s\n", duration, 1e-9 * flop / duration);
|
||||
fflush(stdout);
|
||||
if (NULL != matrix_d) { // Calculate result on the host for validation.
|
||||
dbm_multiply(false, false, 1.0, matrix_a, matrix_b, 1.0, matrix_d, false,
|
||||
1e-8, NULL);
|
||||
|
||||
const double maxeps = 1E-5, epsilon = dbm_maxeps(matrix_d, matrix_c);
|
||||
if (maxeps < epsilon) {
|
||||
printf("ERROR\n");
|
||||
fprintf(stderr, "Failed validation (epsilon=%f).\n", epsilon);
|
||||
exit(1);
|
||||
}
|
||||
} else {
|
||||
printf("ERROR\n");
|
||||
fprintf(stderr, "Failed validation (epsilon=%f).\n", epsilon);
|
||||
exit(1);
|
||||
dbm_release(matrix_d);
|
||||
}
|
||||
|
||||
dbm_mpi_sum_int64(&flop, 1, comm);
|
||||
if (dbm_mpi_comm_rank(comm) == 0) {
|
||||
const double duration = time_end_multiply - time_start_multiply;
|
||||
printf("%6.3f s => %6.1f GFLOP/s\n", duration, 1e-9 * flop / duration);
|
||||
fflush(stdout);
|
||||
}
|
||||
|
||||
dbm_release(matrix_a);
|
||||
dbm_release(matrix_b);
|
||||
dbm_release(matrix_c);
|
||||
dbm_release(matrix_d);
|
||||
}
|
||||
|
||||
/*******************************************************************************
|
||||
|
|
|
|||
|
|
@ -43,7 +43,7 @@ static float *compute_rows_max_eps(const bool trans, const dbm_matrix_t *matrix,
|
|||
const dbm_block_t *blk = &shard->blocks[iblock];
|
||||
const int row = (trans) ? blk->col : blk->row;
|
||||
#pragma omp atomic
|
||||
nblocks_per_row[row]++;
|
||||
++nblocks_per_row[row];
|
||||
}
|
||||
}
|
||||
#pragma omp master
|
||||
|
|
@ -69,6 +69,7 @@ typedef struct {
|
|||
#if defined(__OFFLOAD) && !defined(__NO_OFFLOAD_DBM)
|
||||
dbm_multiply_gpu_context_t gpu;
|
||||
#endif
|
||||
int cpu_options; // Binary or'ed dbm_multiply_cpu_options (enum).
|
||||
} backend_context_t;
|
||||
|
||||
/*******************************************************************************
|
||||
|
|
@ -76,7 +77,9 @@ typedef struct {
|
|||
* \author Ole Schuett
|
||||
******************************************************************************/
|
||||
static backend_context_t *backend_start(const dbm_matrix_t *matrix_c) {
|
||||
backend_context_t *ctx = calloc(1, sizeof(backend_context_t));
|
||||
backend_context_t *const ctx = calloc(1, sizeof(backend_context_t));
|
||||
// BLAS and LIBXSMM benefit in general from DBM_MULTIPLY_TASK_REORDER.
|
||||
ctx->cpu_options = DBM_MULTIPLY_TASK_REORDER;
|
||||
|
||||
#if defined(__OFFLOAD) && !defined(__NO_OFFLOAD_DBM)
|
||||
dbm_multiply_gpu_start(DBM_MAX_BATCH_SIZE, dbm_get_num_shards(matrix_c),
|
||||
|
|
@ -92,16 +95,16 @@ static backend_context_t *backend_start(const dbm_matrix_t *matrix_c) {
|
|||
* \brief Private routine for handing newly arrived packs to the backend.
|
||||
* \author Ole Schuett
|
||||
******************************************************************************/
|
||||
static void backend_upload_packs(const dbm_pack_t *pack_a,
|
||||
static bool backend_upload_packs(const dbm_pack_t *pack_a,
|
||||
const dbm_pack_t *pack_b,
|
||||
backend_context_t *ctx) {
|
||||
|
||||
#if defined(__OFFLOAD) && !defined(__NO_OFFLOAD_DBM)
|
||||
dbm_multiply_gpu_upload_packs(pack_a, pack_b, &ctx->gpu);
|
||||
return dbm_multiply_gpu_upload_packs(pack_a, pack_b, &ctx->gpu);
|
||||
#else
|
||||
(void)pack_a; // mark as used
|
||||
(void)pack_b;
|
||||
(void)ctx;
|
||||
return false;
|
||||
#endif
|
||||
}
|
||||
|
||||
|
|
@ -113,34 +116,26 @@ static void backend_process_batch(const int ntasks,
|
|||
const dbm_task_t batch[ntasks],
|
||||
const double alpha, const dbm_pack_t *pack_a,
|
||||
const dbm_pack_t *pack_b, const int kshard,
|
||||
const bool finish, dbm_shard_t *shard_c,
|
||||
dbm_shard_t *shard_c, const bool finish,
|
||||
const bool force_cpu,
|
||||
backend_context_t *ctx) {
|
||||
// BLAS and LIBXSMM benefit in general from DBM_MULTIPLY_TASK_REORDER.
|
||||
const int cpu_options = DBM_MULTIPLY_TASK_REORDER;
|
||||
|
||||
if (NULL != ctx) {
|
||||
#if defined(__OFFLOAD) && !defined(__NO_OFFLOAD_DBM)
|
||||
dbm_shard_gpu_t *const shard_g = &ctx->gpu.shards_c_dev[kshard];
|
||||
dbm_multiply_gpu_process_batch(ntasks, batch, alpha, kshard, &ctx->gpu);
|
||||
if (finish) { // Start downloading the current shard of matrix_c.
|
||||
// Grow host buffer if necessary.
|
||||
dbm_shard_allocate_promised_blocks(shard_c);
|
||||
// Download results from device.
|
||||
assert(shard_c->data_size == shard_g->data_size);
|
||||
offloadMemcpyAsyncDtoH(shard_c->data, shard_g->data,
|
||||
shard_g->data_size * sizeof(double),
|
||||
shard_g->stream);
|
||||
}
|
||||
#else
|
||||
(void)kshard;
|
||||
(void)finish;
|
||||
dbm_multiply_cpu_process_batch(ntasks, batch, alpha, pack_a, pack_b,
|
||||
shard_c, cpu_options);
|
||||
if (!force_cpu) {
|
||||
dbm_multiply_gpu_process_batch(ntasks, batch, alpha, shard_c, kshard,
|
||||
finish, &ctx->gpu);
|
||||
} else
|
||||
#endif
|
||||
} else { // Validate against host.
|
||||
{
|
||||
(void)kshard;
|
||||
(void)finish;
|
||||
(void)force_cpu;
|
||||
dbm_multiply_cpu_process_batch(ntasks, batch, alpha, pack_a, pack_b,
|
||||
shard_c, ctx->cpu_options);
|
||||
}
|
||||
} else { // Validate against host (aka CPU).
|
||||
dbm_multiply_cpu_process_batch(ntasks, batch, alpha, pack_a, pack_b,
|
||||
shard_c,
|
||||
cpu_options | DBM_MULTIPLY_BLAS_LIBRARY);
|
||||
shard_c, DBM_MULTIPLY_BLAS_LIBRARY);
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -164,9 +159,11 @@ static void multiply_packs(const bool transa, const bool transb,
|
|||
const dbm_pack_t *pack_b,
|
||||
const dbm_matrix_t *matrix_a,
|
||||
const dbm_matrix_t *matrix_b, dbm_matrix_t *matrix_c,
|
||||
const bool retain_sparsity,
|
||||
const float *rows_max_eps, int64_t *flop,
|
||||
backend_context_t *ctx) {
|
||||
const float *rows_max_eps,
|
||||
const bool retain_sparsity, const bool force_cpu,
|
||||
int64_t *flop, backend_context_t *ctx) {
|
||||
// For validation, FLOPS do not count, and relying on ctx is not necessary.
|
||||
backend_context_t *const context = (NULL != flop ? ctx : NULL);
|
||||
const float alpha2 = alpha * alpha;
|
||||
int64_t flop_sum = 0;
|
||||
|
||||
|
|
@ -288,13 +285,13 @@ static void multiply_packs(const bool transa, const bool transb,
|
|||
|
||||
if (ntasks == DBM_MAX_BATCH_SIZE) {
|
||||
backend_process_batch(ntasks, batch, alpha, pack_a, pack_b,
|
||||
ishard, false, shard_c, ctx);
|
||||
ishard, shard_c, false, force_cpu, context);
|
||||
ntasks = 0;
|
||||
}
|
||||
}
|
||||
}
|
||||
backend_process_batch(ntasks, batch, alpha, pack_a, pack_b, ishard,
|
||||
true, shard_c, ctx);
|
||||
shard_c, true, force_cpu, context);
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -322,9 +319,6 @@ void dbm_multiply(const bool transa, const bool transb, const double alpha,
|
|||
assert(omp_get_num_threads() == 1);
|
||||
assert(matrix_a != NULL && matrix_b != NULL && matrix_c != NULL);
|
||||
|
||||
// Compute filter thresholds for each row.
|
||||
float *rows_max_eps = compute_rows_max_eps(transa, matrix_a, filter_eps);
|
||||
|
||||
// Throughout the matrix multiplication code the "sum_index" and "free_index"
|
||||
// denote the summation (aka dummy) and free index from the Einstein notation.
|
||||
const int num_sum_index_a = (transa) ? matrix_a->nrows : matrix_a->ncols;
|
||||
|
|
@ -354,6 +348,9 @@ void dbm_multiply(const bool transa, const bool transb, const double alpha,
|
|||
dbm_copy(matrix_d, matrix_c);
|
||||
}
|
||||
|
||||
// Compute filter thresholds for each row.
|
||||
float *rows_max_eps = compute_rows_max_eps(transa, matrix_a, filter_eps);
|
||||
|
||||
// Start uploading matrix_c to the GPU.
|
||||
backend_context_t *ctx = backend_start(matrix_c);
|
||||
|
||||
|
|
@ -369,9 +366,11 @@ void dbm_multiply(const bool transa, const bool transb, const double alpha,
|
|||
// Main loop.
|
||||
dbm_pack_t *pack_a, *pack_b;
|
||||
while (dbm_comm_iterator_next(iter, &pack_a, &pack_b)) {
|
||||
backend_upload_packs(pack_a, pack_b, ctx);
|
||||
const bool uploaded = backend_upload_packs(pack_a, pack_b, ctx);
|
||||
(void)uploaded; // mark used
|
||||
multiply_packs(transa, transb, alpha, pack_a, pack_b, matrix_a, matrix_b,
|
||||
matrix_c, retain_sparsity, rows_max_eps, flop, ctx);
|
||||
matrix_c, rows_max_eps, retain_sparsity, false /*!uploaded*/,
|
||||
flop, ctx);
|
||||
}
|
||||
|
||||
// Wait for all other MPI ranks to complete, then release ressources.
|
||||
|
|
@ -379,23 +378,19 @@ void dbm_multiply(const bool transa, const bool transb, const double alpha,
|
|||
backend_stop(ctx);
|
||||
|
||||
if (NULL != matrix_d) {
|
||||
int npacks = 0;
|
||||
iter =
|
||||
dbm_comm_iterator_start(transa, transb, matrix_a, matrix_b, matrix_d);
|
||||
while (dbm_comm_iterator_next(iter, &pack_a, &pack_b)) {
|
||||
multiply_packs(transa, transb, alpha, pack_a, pack_b, matrix_a, matrix_b,
|
||||
matrix_d, retain_sparsity, rows_max_eps, NULL, NULL);
|
||||
++npacks;
|
||||
matrix_d, rows_max_eps, retain_sparsity, true, NULL, ctx);
|
||||
}
|
||||
dbm_comm_iterator_stop(iter);
|
||||
const double epsilon = dbm_maxeps(matrix_d, matrix_c);
|
||||
if (maxeps < epsilon) {
|
||||
if (1 == verify) {
|
||||
fprintf(stderr, "WARN ACC/LIBDBM: npacks=%i diff=%g\n", npacks,
|
||||
epsilon);
|
||||
fprintf(stderr, "WARN ACC/LIBDBM: diff=%g\n", epsilon);
|
||||
} else {
|
||||
fprintf(stderr, "ERROR ACC/LIBDBM: npacks=%i diff=%g\n", npacks,
|
||||
epsilon);
|
||||
fprintf(stderr, "ERROR ACC/LIBDBM: diff=%g\n", epsilon);
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -28,9 +28,9 @@ void dbm_multiply_gpu_start(const int max_batch_size, const int nshards,
|
|||
offload_activate_chosen_device();
|
||||
|
||||
ctx->nshards = nshards;
|
||||
ctx->shards_c_host = shards_c_host;
|
||||
ctx->max_batch_size = max_batch_size;
|
||||
offloadStreamCreate(&ctx->main_stream);
|
||||
offloadEventCreate(&ctx->upload_event);
|
||||
|
||||
// Allocate device storage for batches.
|
||||
const size_t size = nshards * max_batch_size * sizeof(dbm_task_t);
|
||||
|
|
@ -40,17 +40,18 @@ void dbm_multiply_gpu_start(const int max_batch_size, const int nshards,
|
|||
ctx->shards_c_dev = malloc(nshards * sizeof(dbm_shard_gpu_t));
|
||||
assert(ctx->shards_c_dev != NULL || nshards == 0);
|
||||
for (int i = 0; i < nshards; i++) {
|
||||
const dbm_shard_t *shard_c_host = &ctx->shards_c_host[i];
|
||||
dbm_shard_gpu_t *shard_c_dev = &ctx->shards_c_dev[i];
|
||||
offloadStreamCreate(&shard_c_dev->stream);
|
||||
shard_c_dev->data_size = shard_c_host->data_size;
|
||||
const dbm_shard_t *const shard_c_host = &shards_c_host[i];
|
||||
dbm_shard_gpu_t *shard_g = &ctx->shards_c_dev[i];
|
||||
offloadStreamCreate(&shard_g->stream);
|
||||
offloadEventCreate(&shard_g->event);
|
||||
shard_g->data_size = shard_c_host->data_size;
|
||||
// only allocate data_size on device rather than data_allocated
|
||||
shard_c_dev->data_allocated = shard_c_host->data_size;
|
||||
shard_c_dev->data =
|
||||
dbm_mempool_device_malloc(shard_c_dev->data_allocated * sizeof(double));
|
||||
offloadMemcpyAsyncHtoD(shard_c_dev->data, shard_c_host->data,
|
||||
shard_c_dev->data_size * sizeof(double),
|
||||
shard_c_dev->stream);
|
||||
shard_g->data_allocated = shard_c_host->data_size;
|
||||
shard_g->data =
|
||||
dbm_mempool_device_malloc(shard_g->data_allocated * sizeof(double));
|
||||
offloadMemcpyAsyncHtoD(shard_g->data, shard_c_host->data,
|
||||
shard_g->data_size * sizeof(double),
|
||||
shard_g->stream);
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -71,29 +72,35 @@ static void upload_pack(const dbm_pack_t *pack_host, dbm_pack_t *pack_dev,
|
|||
|
||||
/*******************************************************************************
|
||||
* \brief Internal routine for uploading newly arrived packs onto the device.
|
||||
* \author Ole Schuett
|
||||
* \author Ole Schuett and Hans Pabst
|
||||
******************************************************************************/
|
||||
void dbm_multiply_gpu_upload_packs(const dbm_pack_t *pack_a,
|
||||
bool dbm_multiply_gpu_upload_packs(const dbm_pack_t *pack_a,
|
||||
const dbm_pack_t *pack_b,
|
||||
dbm_multiply_gpu_context_t *ctx) {
|
||||
// Assume GPU device was activated earlier.
|
||||
// Wait for all c-streams to complete before overwriting old packs.
|
||||
offloadEvent_t event;
|
||||
offloadEventCreate(&event);
|
||||
for (int i = 0; i < ctx->nshards; i++) {
|
||||
offloadEventRecord(event, ctx->shards_c_dev[i].stream);
|
||||
offloadStreamWaitEvent(ctx->main_stream, event);
|
||||
offloadEventRecord(ctx->upload_event, ctx->shards_c_dev[i].stream);
|
||||
offloadStreamWaitEvent(ctx->main_stream, ctx->upload_event);
|
||||
}
|
||||
// Record event to check if all c-streams already completed.
|
||||
offloadEventRecord(ctx->upload_event, ctx->main_stream);
|
||||
|
||||
bool uploaded = false;
|
||||
/*if (offloadEventQuery(ctx->upload_event))*/
|
||||
{
|
||||
upload_pack(pack_a, &ctx->pack_a_dev, ctx->main_stream);
|
||||
upload_pack(pack_b, &ctx->pack_b_dev, ctx->main_stream);
|
||||
|
||||
// Have all c-streams wait until new packs are uploaded.
|
||||
offloadEventRecord(ctx->upload_event, ctx->main_stream);
|
||||
for (int i = 0; i < ctx->nshards; i++) {
|
||||
offloadStreamWaitEvent(ctx->shards_c_dev[i].stream, ctx->upload_event);
|
||||
}
|
||||
uploaded = true;
|
||||
}
|
||||
|
||||
upload_pack(pack_a, &ctx->pack_a_dev, ctx->main_stream);
|
||||
upload_pack(pack_b, &ctx->pack_b_dev, ctx->main_stream);
|
||||
|
||||
// Have all c-streams wait until new packs are uploaded.
|
||||
offloadEventRecord(event, ctx->main_stream);
|
||||
for (int i = 0; i < ctx->nshards; i++) {
|
||||
offloadStreamWaitEvent(ctx->shards_c_dev[i].stream, event);
|
||||
}
|
||||
offloadEventDestroy(event);
|
||||
return uploaded;
|
||||
}
|
||||
|
||||
/*******************************************************************************
|
||||
|
|
@ -101,85 +108,70 @@ void dbm_multiply_gpu_upload_packs(const dbm_pack_t *pack_a,
|
|||
* \author Ole Schuett
|
||||
******************************************************************************/
|
||||
void dbm_multiply_gpu_process_batch(const int ntasks, const dbm_task_t *batch,
|
||||
const double alpha, const int kshard,
|
||||
const double alpha, dbm_shard_t *shard_c,
|
||||
const int kshard, const bool finish,
|
||||
dbm_multiply_gpu_context_t *ctx) {
|
||||
if (ntasks == 0) {
|
||||
return; // Nothing to do.
|
||||
}
|
||||
|
||||
// Assume GPU device was activated earlier.
|
||||
const dbm_shard_t *shard_c_host = &ctx->shards_c_host[kshard];
|
||||
dbm_shard_gpu_t *shard_c_dev = &ctx->shards_c_dev[kshard];
|
||||
assert(NULL != shard_c_host && NULL != shard_c_dev);
|
||||
|
||||
// Upload new batch.
|
||||
dbm_task_t *batch_dev = &ctx->batches_dev[kshard * ctx->max_batch_size];
|
||||
const size_t size = ntasks * sizeof(dbm_task_t);
|
||||
offloadMemcpyAsyncHtoD(batch_dev, batch, size, shard_c_dev->stream);
|
||||
offloadEvent_t memsetup;
|
||||
offloadEventCreate(&memsetup);
|
||||
|
||||
// Reallocate shard_c_dev->data if necessary.
|
||||
dbm_shard_gpu_t *const shard_g = &ctx->shards_c_dev[kshard];
|
||||
double *old_data_dev = NULL;
|
||||
if (shard_c_host->data_promised > shard_c_dev->data_allocated) {
|
||||
shard_c_dev->data_allocated =
|
||||
DBM_OVERCOMMIT_DEVICE * shard_c_host->data_promised;
|
||||
assert(shard_c_host->data_promised <= shard_c_dev->data_allocated);
|
||||
old_data_dev = shard_c_dev->data;
|
||||
shard_c_dev->data =
|
||||
dbm_mempool_device_malloc(shard_c_dev->data_allocated * sizeof(double));
|
||||
// Omit to wait for copy before freeing old buffer.
|
||||
offloadMemcpyAsyncDtoD(shard_c_dev->data, old_data_dev,
|
||||
shard_c_dev->data_size * sizeof(double),
|
||||
shard_c_dev->stream);
|
||||
}
|
||||
offloadEventRecord(memsetup, shard_c_dev->stream);
|
||||
|
||||
// Zero new blocks if necessary.
|
||||
if (shard_c_host->data_promised > shard_c_dev->data_size) {
|
||||
const int tail = shard_c_host->data_promised - shard_c_dev->data_size;
|
||||
offloadMemsetAsync(&shard_c_dev->data[shard_c_dev->data_size], 0,
|
||||
tail * sizeof(double), shard_c_dev->stream);
|
||||
shard_c_dev->data_size = shard_c_host->data_promised;
|
||||
if (0 < ntasks) {
|
||||
assert(NULL != shard_c && NULL != shard_g);
|
||||
|
||||
// Upload new batch.
|
||||
dbm_task_t *batch_dev = &ctx->batches_dev[kshard * ctx->max_batch_size];
|
||||
const size_t size = ntasks * sizeof(dbm_task_t);
|
||||
offloadMemcpyAsyncHtoD(batch_dev, batch, size, shard_g->stream);
|
||||
|
||||
// Reallocate shard_g->data if necessary.
|
||||
if (shard_c->data_promised > shard_g->data_allocated) {
|
||||
shard_g->data_allocated = DBM_OVERCOMMIT_DEVICE * shard_c->data_promised;
|
||||
assert(shard_c->data_promised <= shard_g->data_allocated);
|
||||
old_data_dev = shard_g->data;
|
||||
shard_g->data =
|
||||
dbm_mempool_device_malloc(shard_g->data_allocated * sizeof(double));
|
||||
// Omit to wait for copy before freeing old buffer.
|
||||
offloadMemcpyAsyncDtoD(shard_g->data, old_data_dev,
|
||||
shard_g->data_size * sizeof(double),
|
||||
shard_g->stream);
|
||||
}
|
||||
offloadEventRecord(shard_g->event, shard_g->stream);
|
||||
|
||||
// Zero new blocks if necessary.
|
||||
if (shard_c->data_promised > shard_g->data_size) {
|
||||
const int tail = shard_c->data_promised - shard_g->data_size;
|
||||
offloadMemsetAsync(&shard_g->data[shard_g->data_size], 0,
|
||||
tail * sizeof(double), shard_g->stream);
|
||||
shard_g->data_size = shard_c->data_promised;
|
||||
}
|
||||
|
||||
// Launch kernel.
|
||||
assert(0 != shard_g->data_size);
|
||||
dbm_multiply_gpu_launch_kernel(shard_g->stream, alpha, ntasks, batch,
|
||||
batch_dev, ctx->pack_a_dev.data,
|
||||
ctx->pack_b_dev.data, shard_g->data);
|
||||
OFFLOAD_CHECK(offloadGetLastError());
|
||||
}
|
||||
|
||||
// Launch kernel.
|
||||
assert(0 != shard_c_dev->data_size);
|
||||
dbm_multiply_gpu_launch_kernel(shard_c_dev->stream, alpha, ntasks, batch,
|
||||
batch_dev, ctx->pack_a_dev.data,
|
||||
ctx->pack_b_dev.data, shard_c_dev->data);
|
||||
OFFLOAD_CHECK(offloadGetLastError());
|
||||
|
||||
// Wait for:
|
||||
// - Batch to be uploaded (before refilling it).
|
||||
// - Device memory buffer (if resized).
|
||||
offloadEventSynchronize(memsetup);
|
||||
offloadEventDestroy(memsetup);
|
||||
|
||||
// Safely freeing old buffer.
|
||||
if (NULL != old_data_dev) {
|
||||
dbm_mempool_device_free(old_data_dev);
|
||||
}
|
||||
}
|
||||
|
||||
/*******************************************************************************
|
||||
* \brief Internal routine for downloading results from the device.
|
||||
* \author Ole Schuett
|
||||
******************************************************************************/
|
||||
void dbm_multiply_gpu_download_results(dbm_multiply_gpu_context_t *ctx) {
|
||||
// Assume GPU device was activated earlier.
|
||||
#pragma omp parallel for DBM_OMP_SCHEDULE
|
||||
for (int i = 0; i < ctx->nshards; i++) {
|
||||
if (finish) { // Start downloading the current shard of matrix_c.
|
||||
// Grow host buffer if necessary.
|
||||
dbm_shard_t *shard_c_host = &ctx->shards_c_host[i];
|
||||
dbm_shard_allocate_promised_blocks(shard_c_host);
|
||||
|
||||
dbm_shard_allocate_promised_blocks(shard_c);
|
||||
// Download results from device.
|
||||
dbm_shard_gpu_t *shard_c_dev = &ctx->shards_c_dev[i];
|
||||
assert(shard_c_host->data_size == shard_c_dev->data_size);
|
||||
const size_t size = shard_c_dev->data_size * sizeof(double);
|
||||
offloadMemcpyAsyncDtoH(shard_c_host->data, shard_c_dev->data, size,
|
||||
shard_c_dev->stream);
|
||||
assert(shard_c->data_size == shard_g->data_size);
|
||||
offloadMemcpyAsyncDtoH(shard_c->data, shard_g->data,
|
||||
shard_g->data_size * sizeof(double),
|
||||
shard_g->stream);
|
||||
}
|
||||
|
||||
if (0 < ntasks) {
|
||||
// Wait for:
|
||||
// - Batch to be uploaded (before refilling it).
|
||||
// - Safely freeing device buffer (if resized).
|
||||
offloadEventSynchronize(shard_g->event);
|
||||
|
||||
if (NULL != old_data_dev) {
|
||||
dbm_mempool_device_free(old_data_dev);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -192,10 +184,11 @@ void dbm_multiply_gpu_stop(dbm_multiply_gpu_context_t *ctx) {
|
|||
// Wait for completion, then free gpu ressources.
|
||||
#pragma omp parallel for DBM_OMP_SCHEDULE
|
||||
for (int i = 0; i < ctx->nshards; i++) {
|
||||
dbm_shard_gpu_t *shard_c_dev = &ctx->shards_c_dev[i];
|
||||
offloadStreamSynchronize(shard_c_dev->stream);
|
||||
offloadStreamDestroy(shard_c_dev->stream);
|
||||
dbm_mempool_device_free(shard_c_dev->data);
|
||||
dbm_shard_gpu_t *const shard_g = &ctx->shards_c_dev[i];
|
||||
offloadStreamSynchronize(shard_g->stream);
|
||||
offloadStreamDestroy(shard_g->stream);
|
||||
offloadEventDestroy(shard_g->event);
|
||||
dbm_mempool_device_free(shard_g->data);
|
||||
}
|
||||
free(ctx->shards_c_dev);
|
||||
|
||||
|
|
@ -203,6 +196,7 @@ void dbm_multiply_gpu_stop(dbm_multiply_gpu_context_t *ctx) {
|
|||
dbm_mempool_device_free(ctx->pack_b_dev.data);
|
||||
dbm_mempool_device_free(ctx->batches_dev);
|
||||
offloadStreamDestroy(ctx->main_stream);
|
||||
offloadEventDestroy(ctx->upload_event);
|
||||
}
|
||||
|
||||
#endif // defined(__OFFLOAD) && !defined(__NO_OFFLOAD_DBM)
|
||||
|
|
|
|||
|
|
@ -23,6 +23,7 @@ typedef struct {
|
|||
int data_size;
|
||||
int data_allocated;
|
||||
offloadStream_t stream;
|
||||
offloadEvent_t event;
|
||||
} dbm_shard_gpu_t;
|
||||
|
||||
/*******************************************************************************
|
||||
|
|
@ -31,9 +32,9 @@ typedef struct {
|
|||
******************************************************************************/
|
||||
typedef struct {
|
||||
offloadStream_t main_stream;
|
||||
offloadEvent_t upload_event;
|
||||
|
||||
int nshards;
|
||||
dbm_shard_t *shards_c_host;
|
||||
dbm_shard_gpu_t *shards_c_dev;
|
||||
|
||||
dbm_pack_t pack_a_dev;
|
||||
|
|
@ -53,9 +54,9 @@ void dbm_multiply_gpu_start(const int max_batch_size, const int nshards,
|
|||
|
||||
/*******************************************************************************
|
||||
* \brief Internal routine for uploading newly arrived packs onto the device.
|
||||
* \author Ole Schuett
|
||||
* \author Ole Schuett and Hans Pabst
|
||||
******************************************************************************/
|
||||
void dbm_multiply_gpu_upload_packs(const dbm_pack_t *pack_a,
|
||||
bool dbm_multiply_gpu_upload_packs(const dbm_pack_t *pack_a,
|
||||
const dbm_pack_t *pack_b,
|
||||
dbm_multiply_gpu_context_t *ctx);
|
||||
|
||||
|
|
@ -64,15 +65,10 @@ void dbm_multiply_gpu_upload_packs(const dbm_pack_t *pack_a,
|
|||
* \author Ole Schuett
|
||||
******************************************************************************/
|
||||
void dbm_multiply_gpu_process_batch(const int ntasks, const dbm_task_t *batch,
|
||||
const double alpha, const int kshard,
|
||||
const double alpha, dbm_shard_t *shard_c,
|
||||
const int kshard, const bool finish,
|
||||
dbm_multiply_gpu_context_t *ctx);
|
||||
|
||||
/*******************************************************************************
|
||||
* \brief Internal routine for downloading results from the device.
|
||||
* \author Ole Schuett
|
||||
******************************************************************************/
|
||||
void dbm_multiply_gpu_download_results(dbm_multiply_gpu_context_t *ctx);
|
||||
|
||||
/*******************************************************************************
|
||||
* \brief Internal routine for shutting down the gpu backend.
|
||||
* \author Ole Schuett
|
||||
|
|
|
|||
|
|
@ -22,6 +22,7 @@
|
|||
|
||||
#if defined(__OFFLOAD_CUDA) || defined(__OFFLOAD_HIP) || \
|
||||
defined(__OFFLOAD_OPENCL)
|
||||
#include <stdbool.h>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
|
|
@ -441,6 +442,21 @@ static inline void offloadStreamWaitEvent(offloadStream_t stream,
|
|||
#endif
|
||||
}
|
||||
|
||||
/*******************************************************************************
|
||||
* \brief Wrapper around cudaEventQuery.
|
||||
******************************************************************************/
|
||||
static inline bool offloadEventQuery(offloadEvent_t event) {
|
||||
#if defined(__OFFLOAD_CUDA)
|
||||
return offloadSuccess == cudaEventQuery(event);
|
||||
#elif defined(__OFFLOAD_HIP)
|
||||
return offloadSuccess == hipEventQuery(event);
|
||||
#elif defined(__OFFLOAD_OPENCL)
|
||||
c_dbcsr_acc_bool_t has_occurred;
|
||||
OFFLOAD_CHECK(c_dbcsr_acc_event_query(event, &has_occurred));
|
||||
return (bool)has_occurred;
|
||||
#endif
|
||||
}
|
||||
|
||||
/*******************************************************************************
|
||||
* \brief Wrapper around cudaDeviceSynchronize.
|
||||
******************************************************************************/
|
||||
|
|
|
|||
|
|
@ -24,10 +24,7 @@ FLAG_EXCEPTIONS = (
|
|||
r"__ARM_ARCH",
|
||||
r"__ARM_FEATURE_.+",
|
||||
r"CUDA_VERSION",
|
||||
r"DBM_LIBXSMM_PREFETCH",
|
||||
r"DBM_VALIDATE_AGAINST_.+",
|
||||
r"DBM_ALLOC_.+",
|
||||
r"DBM_MEMPOOL_.+",
|
||||
r"DBM_.+",
|
||||
r"OPENMP_TRACE_SYMBOL",
|
||||
r"OPENCL_.+",
|
||||
r"ACC_OPENCL_.+",
|
||||
|
|
|
|||
Loading…
Add table
Add a link
Reference in a new issue