From 18eff0b89bfedbb6fec7ff61d97dcf215ee6e1dc Mon Sep 17 00:00:00 2001 From: Hans Pabst Date: Tue, 11 Mar 2025 09:17:45 +0100 Subject: [PATCH] DBM: some code cleanup (#4064) - Avoid implicit alloca and prefer (thread-local)malloc. - OpenCL: printing in/homogeneous batches. - Removed tracking extents of task-shapes. - Avoid to compute average flops per rank; apparently done on every call-site. Validation: - Get DBM_MULTIPLY_MAXEPS from environment. - Calibrated default. OpenCL kernel: - Moved BCAST utilities into common header. - Removed BCAST code-path (code reduction). - OpenCL: Adjusted internal parameters - Clear CVEC explicitly. --- src/dbm/dbm_multiply.c | 70 ++++++++++++------------------ src/dbm/dbm_multiply_gpu.c | 5 +-- src/dbm/dbm_multiply_gpu.h | 3 +- src/dbm/dbm_multiply_gpu_kernel.cu | 15 ++++--- src/dbm/dbm_multiply_gpu_kernel.h | 6 +-- src/dbm/dbm_multiply_opencl.c | 68 +++++++++++++++++++---------- src/dbm/dbm_multiply_opencl.cl | 54 ++++++++--------------- 7 files changed, 106 insertions(+), 115 deletions(-) diff --git a/src/dbm/dbm_multiply.c b/src/dbm/dbm_multiply.c index 50c2f69f7b..7e28c3dde7 100644 --- a/src/dbm/dbm_multiply.c +++ b/src/dbm/dbm_multiply.c @@ -7,6 +7,7 @@ #include #include +#include #include #include @@ -27,19 +28,6 @@ #define DBM_VALIDATE_AGAINST_LIBXSMM #endif -/******************************************************************************* - * \brief Updates the min/max of a range of values (initially {INT_MAX, 0}). - * \author Hans Pabst - ******************************************************************************/ -static inline void min_max(int result[2], int value) { - if (value < result[0]) { - result[0] = value; - } - if (result[1] < value) { - result[1] = value; - } -} - /******************************************************************************* * \brief Private routine for computing the max filter threshold for each row. * \author Ole Schuett @@ -89,7 +77,7 @@ typedef struct { } backend_context_t; /******************************************************************************* - * \brief Private routine for intializing the multiplication backend. + * \brief Private routine for initializing the multiplication backend. * \author Ole Schuett ******************************************************************************/ static backend_context_t *backend_start(const dbm_matrix_t *matrix_c) { @@ -127,14 +115,12 @@ static void backend_upload_packs(const dbm_pack_t *pack_a, * \author Ole Schuett ******************************************************************************/ static void backend_process_batch(const int ntasks, dbm_task_t batch[ntasks], - const int mnk_range[3][2], const double alpha, - const dbm_pack_t *pack_a, + const double alpha, const dbm_pack_t *pack_a, const dbm_pack_t *pack_b, const int kshard, dbm_shard_t *shard_c, backend_context_t *ctx) { #if defined(__OFFLOAD) && !defined(__NO_OFFLOAD_DBM) - dbm_multiply_gpu_process_batch(ntasks, batch, mnk_range, alpha, kshard, - &ctx->gpu); + dbm_multiply_gpu_process_batch(ntasks, batch, alpha, kshard, &ctx->gpu); #if defined(DBM_VALIDATE_AGAINST_LIBXSMM) && defined(__LIBXSMM) dbm_shard_gpu_t *const shard_g = &ctx->gpu.shards_c_dev[kshard]; dbm_shard_t shard_r; @@ -162,10 +148,11 @@ static void backend_process_batch(const int ntasks, dbm_task_t batch[ntasks], libxsmm_matdiff_reduce(&diff, &d); } } + const char *const maxeps_env = getenv("DBM_MULTIPLY_MAXEPS"); + const double maxeps = (NULL == maxeps_env ? 1E-13 : fabs(atof(maxeps_env))); const double epsilon = libxsmm_matdiff_epsilon(&diff); - if (1E-15 < epsilon) { - fprintf(stderr, "INFO ACC/LIBDBM: mnk=%ix%ix%i ntasks=%i diff=%g\n", - mnk_range[0][1], mnk_range[1][1], mnk_range[2][1], ntasks, epsilon); + if (maxeps < epsilon) { + fprintf(stderr, "INFO ACC/LIBDBM: diff=%g\n", epsilon); } dbm_shard_release(&shard_r); #else @@ -174,7 +161,6 @@ static void backend_process_batch(const int ntasks, dbm_task_t batch[ntasks], (void)shard_c; // mark as used #endif #else - (void)mnk_range; (void)kshard; (void)ctx; // mark as used dbm_multiply_cpu_process_batch(ntasks, batch, alpha, pack_a, pack_b, shard_c); @@ -221,9 +207,9 @@ static void multiply_packs(const bool transa, const bool transb, const int nshard_rows = matrix_c->dist->rows.nshards; const int nshard_cols = matrix_c->dist->cols.nshards; - int shard_row_start[nshard_rows], shard_col_start[nshard_cols]; - memset(shard_row_start, 0, nshard_rows * sizeof(int)); - memset(shard_col_start, 0, nshard_cols * sizeof(int)); + int *shard_row_start = calloc(nshard_rows, sizeof(int)); + int *shard_col_start = calloc(nshard_cols, sizeof(int)); + assert(NULL != shard_row_start && NULL != shard_col_start); const int *sum_index_sizes_a = (transa) ? matrix_a->row_sizes : matrix_a->col_sizes; @@ -236,6 +222,11 @@ static void multiply_packs(const bool transa, const bool transb, #pragma omp parallel reduction(+ : flop_sum) { + // Thread-private array covering given work in piece-wise fashion. + dbm_task_t *batch = + omp_alloc(sizeof(dbm_task_t) * DBM_MAX_BATCH_SIZE, omp_null_allocator); + assert(NULL != batch); + // Blocks are ordered first by shard. Creating lookup tables of boundaries. #pragma omp for nowait for (int iblock = 1; iblock < pack_a->nblocks; iblock++) { @@ -261,8 +252,6 @@ static void multiply_packs(const bool transa, const bool transb, for (int shard_col = 0; shard_col < nshard_cols; shard_col++) { const int ishard = shard_row * nshard_cols + shard_col; dbm_shard_t *shard_c = &matrix_c->shards[ishard]; - dbm_task_t batch[DBM_MAX_BATCH_SIZE]; - int mnk_range[][2] = {{INT_MAX, 0}, {INT_MAX, 0}, {INT_MAX, 0}}; int ntasks = 0; // Use a merge-join to find pairs of blocks with matching sum indices. @@ -329,27 +318,26 @@ static void multiply_packs(const bool transa, const bool transb, batch[ntasks].offset_a = blk_a->offset; batch[ntasks].offset_b = blk_b->offset; batch[ntasks].offset_c = blk_c->offset; - ntasks++; - - // track MxN-shape covering an entire batch - min_max(mnk_range[0], m); - min_max(mnk_range[1], n); - min_max(mnk_range[2], k); + ++ntasks; if (ntasks == DBM_MAX_BATCH_SIZE) { - backend_process_batch(ntasks, batch, mnk_range, alpha, pack_a, - pack_b, ishard, shard_c, ctx); - mnk_range[0][0] = mnk_range[1][0] = mnk_range[2][0] = INT_MAX; - mnk_range[0][1] = mnk_range[1][1] = mnk_range[2][1] = 0; + backend_process_batch(ntasks, batch, alpha, pack_a, pack_b, + ishard, shard_c, ctx); ntasks = 0; } } } - backend_process_batch(ntasks, batch, mnk_range, alpha, pack_a, pack_b, - ishard, shard_c, ctx); + backend_process_batch(ntasks, batch, alpha, pack_a, pack_b, ishard, + shard_c, ctx); } } + + omp_free(batch, omp_null_allocator); } + + free(shard_row_start); + free(shard_col_start); + *flop += flop_sum; } @@ -408,10 +396,6 @@ void dbm_multiply(const bool transa, const bool transb, const double alpha, free(rows_max_eps); backend_stop(ctx); - // Compute average flops per rank. - dbm_mpi_sum_int64(flop, 1, matrix_c->dist->comm); - *flop = (*flop + matrix_c->dist->nranks - 1) / matrix_c->dist->nranks; - // Final filter pass. dbm_filter(matrix_c, filter_eps); } diff --git a/src/dbm/dbm_multiply_gpu.c b/src/dbm/dbm_multiply_gpu.c index aff48e944a..1579d7dbac 100644 --- a/src/dbm/dbm_multiply_gpu.c +++ b/src/dbm/dbm_multiply_gpu.c @@ -18,7 +18,7 @@ #include /******************************************************************************* - * \brief Internal routine for intializing the gpu backend. + * \brief Internal routine for initializing the gpu backend. * \author Ole Schuett ******************************************************************************/ void dbm_multiply_gpu_start(const int max_batch_size, const int nshards, @@ -103,7 +103,6 @@ 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 int mnk_range[3][2], const double alpha, const int kshard, dbm_multiply_gpu_context_t *ctx) { if (ntasks == 0) { @@ -149,7 +148,7 @@ void dbm_multiply_gpu_process_batch(const int ntasks, const dbm_task_t *batch, // Launch kernel. assert(0 != shard_c_dev->data_size); - dbm_multiply_gpu_launch_kernel(shard_c_dev->stream, mnk_range, alpha, ntasks, + 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()); diff --git a/src/dbm/dbm_multiply_gpu.h b/src/dbm/dbm_multiply_gpu.h index 1a5c3dbe6b..04c78109d5 100644 --- a/src/dbm/dbm_multiply_gpu.h +++ b/src/dbm/dbm_multiply_gpu.h @@ -44,7 +44,7 @@ typedef struct { } dbm_multiply_gpu_context_t; /******************************************************************************* - * \brief Internal routine for intializing the gpu backend. + * \brief Internal routine for initializing the gpu backend. * \author Ole Schuett ******************************************************************************/ void dbm_multiply_gpu_start(const int max_batch_size, const int nshards, @@ -64,7 +64,6 @@ 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 int mnk_range[3][2], const double alpha, const int kshard, dbm_multiply_gpu_context_t *ctx); diff --git a/src/dbm/dbm_multiply_gpu_kernel.cu b/src/dbm/dbm_multiply_gpu_kernel.cu index 51083d552e..43f8d16e36 100644 --- a/src/dbm/dbm_multiply_gpu_kernel.cu +++ b/src/dbm/dbm_multiply_gpu_kernel.cu @@ -192,16 +192,19 @@ __global__ static void process_batch_kernel(const double alpha, * All arguments are assumed to be device pointers. * \author Ole Schuett ******************************************************************************/ -void dbm_multiply_gpu_launch_kernel( - const offloadStream_t stream, const int mnk_range[3][2], const double alpha, - const int ntasks, const dbm_task_t *batch, const double *pack_a_data, - const double *pack_b_data, double *shard_c_data) { +void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, + const double alpha, const int ntasks, + const dbm_task_t *tasks_host, + const dbm_task_t *tasks, + const double *pack_a_data, + const double *pack_b_data, + double *shard_c_data) { const int nblocks = ntasks; // TODO tune launch parameters. const int threads_per_block = NUM_THREADS; const size_t smem_per_block = 0; - (void)mnk_range; // mark used + (void)tasks_host; // mark used process_batch_kernel<<>>( - alpha, batch, pack_a_data, pack_b_data, shard_c_data); + alpha, tasks, pack_a_data, pack_b_data, shard_c_data); } #endif // defined(__OFFLOAD) && !defined(__NO_OFFLOAD_DBM) diff --git a/src/dbm/dbm_multiply_gpu_kernel.h b/src/dbm/dbm_multiply_gpu_kernel.h index 561f4e18bf..7c74561cba 100644 --- a/src/dbm/dbm_multiply_gpu_kernel.h +++ b/src/dbm/dbm_multiply_gpu_kernel.h @@ -23,9 +23,9 @@ extern "C" { * \author Ole Schuett ******************************************************************************/ void dbm_multiply_gpu_launch_kernel( - const offloadStream_t stream, const int mnk_range[3][2], const double alpha, - const int ntasks, const dbm_task_t *batch, const double *pack_a_data, - const double *pack_b_data, double *shard_c_data); + const offloadStream_t stream, const double alpha, const int ntasks, + const dbm_task_t *tasks_host, const dbm_task_t *tasks, + const double *pack_a_data, const double *pack_b_data, double *shard_c_data); #ifdef __cplusplus } diff --git a/src/dbm/dbm_multiply_opencl.c b/src/dbm/dbm_multiply_opencl.c index 29cd579ed4..5d76851137 100644 --- a/src/dbm/dbm_multiply_opencl.c +++ b/src/dbm/dbm_multiply_opencl.c @@ -11,9 +11,32 @@ #include "dbm_multiply_gpu_kernel.h" #include "dbm_multiply_opencl.cl.h" -void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, - const int mnk[3][2], double alpha, - int ntasks, const dbm_task_t *tasks, +typedef struct { + int max_m, avg_m, avg_n, avg_k, changes; +} dbm_multiply_gpu_launch_info_t; + +static void dbm_multiply_gpu_launch_info(dbm_multiply_gpu_launch_info_t *info, + const dbm_task_t *tasks, int ntasks) { + int i = 1; + info->max_m = tasks[0].m; + info->avg_m = tasks[0].m; + info->avg_n = tasks[0].n; + info->avg_k = tasks[0].k; + for (info->changes = 0; i < ntasks; ++i) { + const int m = tasks[i].m, n = tasks[i].n, k = tasks[i].k; + info->max_m = imax(info->max_m, m); + if (info->avg_m != m || info->avg_n != n || info->avg_k != k) { + info->avg_m = (info->avg_m + m) / 2; + info->avg_n = (info->avg_n + n) / 2; + info->avg_k = (info->avg_k + k) / 2; + ++info->changes; + } + } +} + +void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, double alpha, + int ntasks, const dbm_task_t *tasks_host, + const dbm_task_t *tasks, const double *pack_a_data, const double *pack_b_data, double *shard_c_data) { @@ -27,14 +50,12 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, cl_event event, *const perf_event = ((0 <= verbosity && 2 >= verbosity) ? NULL : &event); const c_dbcsr_acc_opencl_stream_t *const str = ACC_OPENCL_STREAM(stream); - const size_t max_m = mnk[0][1], work_tasks = ntasks; - size_t work_size[] = {1, 1, 1}, ibatch = 0, iadata = 0, ibdata = 0, - icdata = 0; + size_t work_size[] = {1, 1, 1}, ibatch = 0; + size_t iadata = 0, ibdata = 0, icdata = 0; + const size_t work_tasks = ntasks; + dbm_multiply_gpu_launch_info_t info = {0}; c_dbcsr_acc_opencl_info_memptr_t adata, bdata, cdata, batch; assert(NULL != pack_a_data && NULL != pack_b_data && NULL != shard_c_data); - assert(0 < mnk[0][0] && 0 < mnk[0][1] && mnk[0][0] <= mnk[0][1]); - assert(0 < mnk[1][0] && 0 < mnk[1][1] && mnk[1][0] <= mnk[1][1]); - assert(0 < mnk[2][0] && 0 < mnk[2][1] && mnk[2][0] <= mnk[2][1]); assert(NULL != str && NULL != str->queue); assert(0 < ntasks && NULL != tasks); #if defined(OPENCL_DBM_SOURCE_MULTIPLY) @@ -49,9 +70,7 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, const char *const bn_env = getenv("DBM_MULTIPLY_BN"); const char *const sm_env = getenv("DBM_MULTIPLY_SM"); const char *const wg_env = getenv("DBM_MULTIPLY_WG"); - const int bn0 = (0 == c_dbcsr_acc_opencl_config.device.nv - ? (0 == c_dbcsr_acc_opencl_config.device.amd ? 4 : 8) - : 2); + const int bn0 = (0 == c_dbcsr_acc_opencl_config.device.nv ? 8 : 4); int lu = LIBXSMM_CLMP(NULL == lu_env ? 0 : atoi(lu_env), -2, 1); int sm = (NULL == sm_env ? 0 /*default*/ : atoi(sm_env)); int bn = LIBXSMM_CLMP( @@ -185,12 +204,13 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 8, cdata.memory); result |= clSetKernelArg(kernel, 9, sizeof(cl_uint), &zero /*C_shape0*/); } else { + size_t size = work_tasks; + dbm_multiply_gpu_launch_info(&info, tasks_host, ntasks); + size *= info.max_m; + /* fixup to be a multiple of the WG-size */ + work_size[0] = (0 < wgsize[0] ? LIBXSMM_UP(size, wgsize[0]) : size); result |= clSetKernelArg(kernel, 2, sizeof(cl_int), &ntasks); - work_size[0] = work_tasks * max_m; - result |= clSetKernelArg(kernel, 3, sizeof(cl_int), work_size); - if (0 < wgsize[0]) { /* fixup to be a multiple of the WG-size */ - work_size[0] = LIBXSMM_UP(work_size[0], wgsize[0]); - } + result |= clSetKernelArg(kernel, 3, sizeof(cl_int), &size); result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 4, batch.memory); result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 5, adata.memory); result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 6, bdata.memory); @@ -209,14 +229,18 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, EXIT_SUCCESS == clGetEventProfilingInfo(*perf_event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL)) { - const double flops = 1E-6 * mnk[0][1] * mnk[1][1] * mnk[2][1] * ntasks; const double dkrnl = 1E-9 * LIBXSMM_DELTA(begin, end); const double dtotl = 1E+3 * LIBXSMM_MAX(dkrnl, dhost); + int pure; + if (1 < ndims) { /* DBM_MULTIPLY_GEN */ + dbm_multiply_gpu_launch_info(&info, tasks_host, ntasks); + } + pure = (100 * (ntasks - info.changes) + ntasks - 1) / ntasks; fprintf(stderr, - "INFO ACC/LIBDBM: DBM-kernel mnk=%ix%ix%i ntasks=%i " - "kernel_ms=%.2g total_ms=%.2g gflops=%.1f\n", - mnk[0][1], mnk[1][1], mnk[2][1], ntasks, dkrnl, dtotl, - flops / dtotl); + "INFO ACC/LIBDBM: DBM-kernel mnk=%ix%ix%i pure=%i%% " + "ntasks=%i kernel_ms=%.2g total_ms=%.2g gflops=%.1f\n", + info.avg_m, info.avg_n, info.avg_k, pure, ntasks, dkrnl, dtotl, + 2E-6 * info.avg_m * info.avg_n * info.avg_k * ntasks / dtotl); } } OFFLOAD_CHECK(result); diff --git a/src/dbm/dbm_multiply_opencl.cl b/src/dbm/dbm_multiply_opencl.cl index 366541b6b5..c87a86c760 100644 --- a/src/dbm/dbm_multiply_opencl.cl +++ b/src/dbm/dbm_multiply_opencl.cl @@ -10,15 +10,6 @@ #include "../../exts/dbcsr/src/acc/opencl/common/opencl_atomics.h" #include "dbm_internal.h" -#if defined(WG) && (0 < WG) && defined(GPU) && (200 <= ACC_OPENCL_VERSION) -#if defined(SG) && (0 < SG) -#define BCST_WG(V) sub_group_broadcast(V, 0) -#else -#define BCST_WG(V) work_group_broadcast(V, 0) -#endif -#endif -#define BCST_NO(V) (V) - #define SINT short #define X(T, I) (T)->I /* task can be taken by value or by pointer */ @@ -32,34 +23,32 @@ #define DBM_MULTIPLY_STORE(ALPHA, TASK, CMAT, CVEC, M, N0, N1) \ do { /* CMAT atomically accumulates CVEC */ \ UNROLL_AUTO for (SINT n = 0; n < (N1); ++n) { /* flush to global */ \ - const int idx = IDT(M, n + (N0), XM(TASK), XN(TASK)); \ - ACCUMULATE((CMAT) + XC(TASK) + idx, (ALPHA) * (CVEC)[n]); \ - (CVEC)[n] = ZERO; /* reset */ \ + const int idx = IDT(M, n + (N0), XM(TASK), XN(TASK)) + XC(TASK); \ + ACCUMULATE((CMAT) + idx, (ALPHA) * (CVEC)[n]); \ } \ } while (0) -#define DBM_MULTIPLY_KERNEL(TASK, AMAT, BMAT, CVEC, M, N0, N1, BCST) \ +#define DBM_MULTIPLY_KERNEL(TASK, AMAT, BMAT, CVEC, M, N0, BN) \ do { /* CVEC accumulates result */ \ UNROLL_AUTO for (SINT k = 0; k < XK(TASK); ++k) { \ const double a = (AMAT)[XA(TASK) + IDT(M, k, XM(TASK), XK(TASK))]; \ - UNROLL_AUTO for (SINT n = 0; n < (N1); ++n) { \ - const int idx = IDX(k, n + (N0), XK(TASK), XN(TASK)); \ - (CVEC)[n] = MAD(a, BCST((BMAT)[idx]), (CVEC)[n]); \ + const int idx = IDX(k, N0, XK(TASK), XN(TASK)); \ + UNROLL_AUTO for (SINT n = 0; n < (BN); ++n) { \ + (CVEC)[n] = MAD(a, (BMAT)[idx + n], (CVEC)[n]); \ } \ } \ } while (0) -#define DBM_MULTIPLY(ALPHA, TASK, AMAT, BMAT, CMAT, CVEC, M, BN, BCST) \ +#define DBM_MULTIPLY(ALPHA, TASK, AMAT, BMAT, CMAT, CVEC, M, BN) \ do { /* DBM_MULTIPLY_KERNEL specialized over N */ \ - SINT n0 = 0; \ - UNROLL_AUTO for (SINT n = 0; n < (BN); ++n) { (CVEC)[n] = ZERO; } \ - if ((BN) <= XN(TASK)) { \ - UNROLL_OUTER(1) for (; (n0 + (BN)) <= XN(TASK); n0 += (BN)) { \ - DBM_MULTIPLY_KERNEL(TASK, AMAT, BMAT, CVEC, M, n0, BN, BCST); \ - DBM_MULTIPLY_STORE(ALPHA, TASK, CMAT, CVEC, M, n0, BN); \ - } \ + SINT n0 = 0, n1 = XN(TASK) - (BN); \ + UNROLL_FORCE(BN) for (SINT n = 0; n < (BN); ++n) { (CVEC)[n] = ZERO; } \ + UNROLL_OUTER(1) for (; n0 <= n1; n0 += (BN)) { \ + DBM_MULTIPLY_KERNEL(TASK, AMAT, BMAT, CVEC, M, n0, BN); \ + DBM_MULTIPLY_STORE(ALPHA, TASK, CMAT, CVEC, M, n0, BN); \ + UNROLL_FORCE(BN) for (SINT n = 0; n < (BN); ++n) { (CVEC)[n] = ZERO; } \ } \ - DBM_MULTIPLY_KERNEL(TASK, AMAT, BMAT, CVEC, M, n0, XN(TASK) - n0, BCST); \ + DBM_MULTIPLY_KERNEL(TASK, AMAT, BMAT, CVEC, M, n0, XN(TASK) - n0); \ DBM_MULTIPLY_STORE(ALPHA, TASK, CMAT, CVEC, M, n0, XN(TASK) - n0); \ } while (0) @@ -79,26 +68,19 @@ dbm_multiply(double alpha, int itask, int ntasks, int size, #else double cvec[BN]; #endif -#if defined(BCST_WG) +#if defined(WG) && (0 < WG) && !defined(NDEBUG) if (i < size) #endif - { + { /* valid task */ const int max_m = size / ntasks, tid = i / max_m; const SINT m = i - tid * max_m; global const dbm_task_t *const task = &tasks[itask + tid]; #if !defined(NDEBUG) if (m < XM(task)) #endif - { /* valid task */ + { /* valid slice (subtask) */ bmat += XB(task); -#if defined(BCST_WG) - if (XM(task) <= XN(task)) { /* BCST_WG to broadcast B-values */ - DBM_MULTIPLY(alpha, task, amat, bmat, cmat, cvec, m, BN, BCST_WG); - } else -#endif - { - DBM_MULTIPLY(alpha, task, amat, bmat, cmat, cvec, m, BN, BCST_NO); - } + DBM_MULTIPLY(alpha, task, amat, bmat, cmat, cvec, m, BN); } } }