mirror of
https://github.com/cp2k/cp2k.git
synced 2026-07-27 13:45:19 -04:00
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.
This commit is contained in:
Hans Pabst
GitHub
parent
2431d80646
commit
18eff0b89b
7 changed files with 106 additions and 115 deletions
|
|
@ -7,6 +7,7 @@
|
|||
|
||||
#include <assert.h>
|
||||
#include <limits.h>
|
||||
#include <omp.h>
|
||||
#include <stdlib.h>
|
||||
#include <string.h>
|
||||
|
||||
|
|
@ -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);
|
||||
}
|
||||
|
|
|
|||
|
|
@ -18,7 +18,7 @@
|
|||
#include <stdio.h>
|
||||
|
||||
/*******************************************************************************
|
||||
* \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());
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
|
||||
|
|
|
|||
|
|
@ -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<<<nblocks, threads_per_block, smem_per_block, stream>>>(
|
||||
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)
|
||||
|
|
|
|||
|
|
@ -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
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
|||
Loading…
Add table
Add a link
Reference in a new issue