From 7face5689eecea11e162328edf493523d06fda3a Mon Sep 17 00:00:00 2001 From: Hans Pabst Date: Thu, 4 Sep 2025 12:59:12 +0200 Subject: [PATCH] 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). --- src/dbm/dbm_miniapp.c | 47 +++--- src/dbm/dbm_multiply.c | 85 +++++----- src/dbm/dbm_multiply_gpu.c | 196 +++++++++++------------ src/dbm/dbm_multiply_gpu.h | 16 +- src/offload/offload_runtime.h | 16 ++ tools/precommit/check_file_properties.py | 5 +- 6 files changed, 184 insertions(+), 181 deletions(-) diff --git a/src/dbm/dbm_miniapp.c b/src/dbm/dbm_miniapp.c index 22eb2b3a65..b6c318364f 100644 --- a/src/dbm/dbm_miniapp.c +++ b/src/dbm/dbm_miniapp.c @@ -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); } /******************************************************************************* diff --git a/src/dbm/dbm_multiply.c b/src/dbm/dbm_multiply.c index e6a7c3c200..a0fef305d0 100644 --- a/src/dbm/dbm_multiply.c +++ b/src/dbm/dbm_multiply.c @@ -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); } } diff --git a/src/dbm/dbm_multiply_gpu.c b/src/dbm/dbm_multiply_gpu.c index edae398117..89d58e6c52 100644 --- a/src/dbm/dbm_multiply_gpu.c +++ b/src/dbm/dbm_multiply_gpu.c @@ -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) diff --git a/src/dbm/dbm_multiply_gpu.h b/src/dbm/dbm_multiply_gpu.h index 04c78109d5..773a0b11ab 100644 --- a/src/dbm/dbm_multiply_gpu.h +++ b/src/dbm/dbm_multiply_gpu.h @@ -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 diff --git a/src/offload/offload_runtime.h b/src/offload/offload_runtime.h index a345a03ffb..bacb18f091 100644 --- a/src/offload/offload_runtime.h +++ b/src/offload/offload_runtime.h @@ -22,6 +22,7 @@ #if defined(__OFFLOAD_CUDA) || defined(__OFFLOAD_HIP) || \ defined(__OFFLOAD_OPENCL) +#include #include #include @@ -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. ******************************************************************************/ diff --git a/tools/precommit/check_file_properties.py b/tools/precommit/check_file_properties.py index a5ca465f99..8114ca4e24 100755 --- a/tools/precommit/check_file_properties.py +++ b/tools/precommit/check_file_properties.py @@ -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_.+",