comp-chem/cp2k
1
0
Fork 0
mirror of https://github.com/cp2k/cp2k.git synced 2026-07-27 21:55:16 -04:00
Code Issues Projects Releases Packages Wiki Activity Actions

DBM/OpenCL: DBCSR/DBM kernel interoperability (#4126)

Browse source 
- Revised parameter format description (support DBCSR-style in kernel).
- DBM_MULTIPLY_SMM=-1: use DBM-kernels for DBCSR.
- DBM_MULTIPLY_SMM=1: use DBCSR-kernels for DBM.

Other:
- Bulk-Lock c_dbcsr_acc_opencl_info_devptr (c_dbcsr_acc_opencl_info_devptr_lock).
- Call dbm_multiply_gpu_launch_info only if necessary.
- Pass stream in non-const fashion.
- Miniapp: drop CFLAGS (linkage).
This commit is contained in:
Hans Pabst 2025-04-08 10:37:23 +02:00 • committed by GitHub
parent f2d5b0afe4
commit b824476083
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
5 changed files with 207 additions and 135 deletions
Download patch file Download diff file Expand all files Collapse all files

6
src/dbm/Makefile
View file

@ -45,8 +45,8 @@ CFLAGS += -fopenmp -Wno-vla-parameter
CC := $(if $(filter-out 0,$(MPI)),mpicc,gcc)
else
MKL_FCRTL := intel
LIBS += $(if $(OMPRT),-l$(OMPRT),-qopenmp)
CFLAGS += -qopenmp
LIBS += $(if $(OMPRT),-fopenmp -Wno-recommended-option,-qopenmp)
CFLAGS += $(if $(OMPRT),-fopenmp -Wno-recommended-option,-qopenmp)
CC := $(if $(filter-out 0,$(MPI)),mpiicx,icx)
endif
@ -206,6 +206,6 @@ endif
cd $(dir $<); $(CC) -c -std=c11 $(CFLAGS) $(notdir $<)
dbm_miniapp.x: dbm_miniapp.o $(ALL_OBJECTS)
$(CC) $(CFLAGS) -o $@ $^ $(LIBS)
$(CC) -o $@ $^ $(LIBS)
#EOF

5
src/dbm/dbm_multiply_gpu_kernel.cu
View file

@ -192,9 +192,8 @@ __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 double alpha, const int ntasks,
const dbm_task_t *tasks_host,
void dbm_multiply_gpu_launch_kernel(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,

10
src/dbm/dbm_multiply_gpu_kernel.h
View file

@ -22,10 +22,12 @@ extern "C" {
* All arguments are assumed to be device pointers.
* \author Ole Schuett
******************************************************************************/
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);
void dbm_multiply_gpu_launch_kernel(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);
#ifdef __cplusplus
}

217
src/dbm/dbm_multiply_opencl.c
View file

@ -19,53 +19,102 @@
#define DBM_TIMER_TICK() libxsmm_timer_tick()
#define DBM_TIMER_TICKINT libxsmm_timer_tickint
int dbm_multiply_opencl_launch_kernel(void *stream, double alpha, int ntasks,
int param_format, const int *params_host,
const int *params,
const double *pack_a_data,
const double *pack_b_data,
double *shard_c_data);
#if defined(OPENCL_LIBSMM_PFORMAT) && (0 < OPENCL_LIBSMM_PFORMAT)
int dbm_multiply_opencl_initialized /*= 0*/;
int dbm_multiply_opencl_smm /*= 0*/;
LIBXSMM_ATTRIBUTE_CTOR static void dbm_multiply_opencl_initialize(void) {
const char *const smm_env = getenv("DBM_MULTIPLY_SMM");
const int smm = (NULL == smm_env ? 0 /*default*/ : atoi(smm_env));
dbm_multiply_opencl_smm =
LIBXSMM_MIN(1 != smm ? smm : 64, (1 << (OPENCL_LIBSMM_PFORMAT - 1)) - 1);
if (0 > dbm_multiply_opencl_smm) {
opencl_libsmm_acc_set_dbm_launch_fn(dbm_multiply_opencl_launch_kernel);
}
++dbm_multiply_opencl_initialized;
}
#endif
typedef struct {
int max_m, max_n, max_k, mnk_count;
int max_m, max_n, max_k, mnk_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 avg_m = tasks[0].m, avg_n = tasks[0].n, avg_k = tasks[0].k, i = 1;
info->max_m = avg_m;
info->max_n = avg_n;
info->max_k = avg_k;
for (info->mnk_count = 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);
info->max_n = imax(info->max_n, n);
info->max_k = imax(info->max_k, k);
if (m != avg_m || n != avg_n || k != avg_k) {
avg_m = (avg_m + m) / 2;
avg_n = (avg_n + n) / 2;
avg_k = (avg_k + k) / 2;
++info->mnk_count;
const int *params, int ntasks,
int param_format) {
if (0 == param_format) { /* native */
const int stride = sizeof(dbm_task_t) / sizeof(int);
int avg_m = params[0], avg_n = params[1], avg_k = params[2], i = stride;
info->max_m = avg_m;
info->max_n = avg_n;
info->max_k = avg_k;
for (info->mnk_changes = 0; i < (ntasks * stride); i += stride) {
const int m = params[i + 0], n = params[i + 1], k = params[i + 2];
info->max_m = imax(info->max_m, m);
info->max_n = imax(info->max_n, n);
info->max_k = imax(info->max_k, k);
if (m != avg_m || n != avg_n || k != avg_k) { /* approximation */
avg_m = (avg_m + m) / 2;
avg_n = (avg_n + n) / 2;
avg_k = (avg_k + k) / 2;
++info->mnk_changes;
}
}
} else {
#if defined(OPENCL_LIBSMM_PFORMAT) && (0 < OPENCL_LIBSMM_PFORMAT)
const int mask = (1 << OPENCL_LIBSMM_PFORMAT) - 1;
info->max_m = mask & (param_format);
info->max_n = mask & (param_format >> (OPENCL_LIBSMM_PFORMAT));
info->max_k = mask & (param_format >> (OPENCL_LIBSMM_PFORMAT * 2));
info->mnk_changes = 0; /* homogeneous */
#else
assert(0);
#endif
}
}
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) {
static void dbm_multiply_opencl_print(FILE *stream, const char *name, int val) {
if (0 != val) {
fprintf(stream, " %s=%i", name, val);
}
}
int dbm_multiply_opencl_launch_kernel(void *stream, double alpha, int ntasks,
int param_format, const int *params_host,
const int *params,
const double *pack_a_data,
const double *pack_b_data,
double *shard_c_data) {
const DBM_TIMER_TICKINT start = DBM_TIMER_TICK();
const c_dbcsr_acc_opencl_config_t *const config = &c_dbcsr_acc_opencl_config;
#if defined(OPENCL_LIBSMM_PFORMAT)
const char *const smm_env = getenv("DBM_MULTIPLY_SMM");
int max_kernel_dim = (NULL == smm_env ? 0 /*default*/ : atoi(smm_env));
#endif
const int verbosity = config->verbosity;
int result = EXIT_SUCCESS;
cl_event e = NULL, *const event =
((0 <= verbosity && 2 >= verbosity) ? NULL : &e);
dbm_multiply_gpu_launch_info_t info = {0};
dbm_multiply_gpu_launch_info(&info, tasks_host, ntasks);
assert(NULL != pack_a_data && NULL != pack_b_data && NULL != shard_c_data);
assert(0 < ntasks && NULL != tasks);
#if defined(OPENCL_LIBSMM_PFORMAT)
if (0 != info.mnk_count || 1 != alpha ||
(max_kernel_dim * max_kernel_dim) < (info.max_m * info.max_n))
assert(NULL != params_host || 0 == ntasks);
assert(NULL != params || 0 == ntasks);
if (0 == ntasks) {
return result;
}
#if defined(OPENCL_LIBSMM_PFORMAT) && (0 < OPENCL_LIBSMM_PFORMAT)
if (0 == dbm_multiply_opencl_initialized) {
dbm_multiply_opencl_initialize();
}
if (0 != dbm_multiply_opencl_smm) {
dbm_multiply_gpu_launch_info(&info, params_host, ntasks, param_format);
}
if (0 > dbm_multiply_opencl_smm || dbm_multiply_opencl_smm < info.max_m ||
dbm_multiply_opencl_smm < info.max_n ||
dbm_multiply_opencl_smm < info.max_k || 0 == info.max_k || 1 != alpha)
#endif
{
#if defined(OPENCL_DBM_SOURCE_MULTIPLY)
@ -75,7 +124,12 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, double alpha,
static int ndims = 1, clinear = 0;
static size_t wgsize[] = {0, 0, 0};
const c_dbcsr_acc_opencl_stream_t *const str = ACC_OPENCL_STREAM(stream);
const c_dbcsr_acc_opencl_device_t *const devinfo = &config->device;
ACC_OPENCL_LOCKTYPE *const lock_memory =
(NULL != devinfo->clSetKernelArgMemPointerINTEL ? NULL
: config->lock_memory);
c_dbcsr_acc_opencl_info_memptr_t adata, bdata, cdata, batch;
const int stride = (0 == param_format ? 6 : 3);
size_t work_size[] = {1, 1, 1}, ibatch = 0;
size_t iadata = 0, ibdata = 0, icdata = 0;
const size_t work_tasks = ntasks;
@ -83,7 +137,7 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, double alpha,
if (NULL == kernel_global) { /* initial check if kernel is present */
ACC_OPENCL_ACQUIRE(config->lock_main);
if (NULL == kernel_global) {
char params[ACC_OPENCL_BUFFERSIZE] =
char flags[ACC_OPENCL_BUFFERSIZE] =
"-cl-fast-relaxed-math -cl-denorms-are-zero";
const char *const gen_env = getenv("DBM_MULTIPLY_GEN");
const char *const lin_env = getenv("DBM_MULTIPLY_LIN");
@ -92,33 +146,32 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, double alpha,
const char *const wg_env = getenv("DBM_MULTIPLY_WG");
const char *const lu_env = getenv("DBM_MULTIPLY_LU");
const char *const xf_env = getenv("DBM_MULTIPLY_XF");
const c_dbcsr_acc_opencl_device_t *const devinfo = &config->device;
int sm = (NULL == sm_env ? 0 /*default*/ : atoi(sm_env));
const int bn0 = (0 == devinfo->nv ? (0 == devinfo->amd ? 4 : 8) : 2);
const int bn1 = ((0 == sm && 0 == clinear) ? bn0 : (bn0 * 2));
int bn = LIBXSMM_CLMP(NULL == bn_env ? bn1 : atoi(bn_env), 1, 32);
int lu = LIBXSMM_CLMP(NULL == lu_env ? 0 : atoi(lu_env), -2, 1);
int gen = ((NULL == bn_env && NULL == sm_env && NULL == wg_env &&
NULL == lu_env && NULL == lin_env)
NULL == lu_env && NULL == lin_env && 0 == param_format)
? (NULL == gen_env ? 1 /*default*/ : atoi(gen_env))
: 0);
const int gpu = (CL_DEVICE_TYPE_GPU == devinfo->type);
const int xf = (NULL == xf_env ? -1 /*default*/ : atoi(xf_env));
const char *extensions[] = {NULL, NULL}, *flags = NULL;
const char *extensions[] = {NULL, NULL}, *options = NULL;
size_t nextensions = sizeof(extensions) / sizeof(*extensions);
const size_t wgsize0 = devinfo->wgsize[0], wgsize1 = devinfo->wgsize[1];
size_t wgsize2 = devinfo->wgsize[2];
size_t offset =
((0 == config->debug && 0 == config->dump) ? strlen(params) : 0);
((0 == config->debug && 0 == config->dump) ? strlen(flags) : 0);
offset += (size_t)c_dbcsr_acc_opencl_flags_atomics(
devinfo, c_dbcsr_acc_opencl_atomic_fp_64, extensions, &nextensions,
params + offset, sizeof(params) - offset);
flags + offset, sizeof(flags) - offset);
if (2 <= gen ||
(0 != gen && 0 != wgsize2 /*subgroups*/ &&
2 <= *devinfo->std_level && NULL != extensions[1] &&
NULL != strstr(extensions[1], "cl_ext_float_atomics"))) {
offset +=
(size_t)LIBXSMM_SNPRINTF(params + offset, sizeof(params) - offset,
(size_t)LIBXSMM_SNPRINTF(flags + offset, sizeof(flags) - offset,
" -DDBM_MULTIPLY_OPENCL_GEN");
wgsize[1] = wgsize[2] = 1;
wgsize[0] = 16;
@ -143,45 +196,31 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, double alpha,
: 0);
clinear = (NULL == lin_env ? 0 /*default*/ : atoi(lin_env));
offset += (size_t)LIBXSMM_SNPRINTF(
params + offset, sizeof(params) - offset,
flags + offset, sizeof(flags) - offset,
" %s %s -DBN=%i -DSM=%i -DLU=%i -DWG=%i -DSG=%i",
0 != gpu ? "-DGPU" : "", 0 == clinear ? "" : "-DCLINEAR", bn, sm,
lu, (int)wgsize[0], (int)wgsize2);
gen = 0;
}
if (0 != devinfo->intel && 0 < xf) {
flags = "-cl-intel-256-GRF-per-thread";
options = "-cl-intel-256-GRF-per-thread";
}
result |= (sizeof(params) > offset ? EXIT_SUCCESS : EXIT_FAILURE);
result |= (sizeof(flags) > offset ? EXIT_SUCCESS : EXIT_FAILURE);
result |= c_dbcsr_acc_opencl_kernel(
0 /*source_is_file*/, OPENCL_DBM_SOURCE_MULTIPLY, "dbm_multiply",
params, flags, NULL /*try*/, NULL /*try_ok*/, extensions,
flags, options, NULL /*try*/, NULL /*try_ok*/, extensions,
nextensions, &kernel_global);
if (2 <= verbosity || 0 > verbosity) {
if (EXIT_SUCCESS == result) {
const double ds = DBM_TIMER_DIFF(start, DBM_TIMER_TICK());
fprintf(stderr, "INFO ACC/LIBDBM: DBM-kernel gpu=%i", gpu);
if (0 != gen) { /* generated kernel */
fprintf(stderr, " gen=%i", gen);
}
if (0 != clinear) {
fprintf(stderr, " lin=%i", clinear);
}
if (0 != bn) {
fprintf(stderr, " bn=%i", bn);
}
if (0 != sm) {
fprintf(stderr, " sm=%i", sm);
}
if (0 != wgsize[0]) {
fprintf(stderr, " wg=%i", (int)wgsize[0]);
}
if (0 != wgsize2) {
fprintf(stderr, " sg=%i", (int)wgsize2);
}
if (0 != lu) {
fprintf(stderr, " lu=%i", lu);
}
dbm_multiply_opencl_print(stderr, "gen", gen); /* generated */
dbm_multiply_opencl_print(stderr, "lin", clinear);
dbm_multiply_opencl_print(stderr, "bn", bn);
dbm_multiply_opencl_print(stderr, "sm", sm);
dbm_multiply_opencl_print(stderr, "wg", (int)wgsize[0]);
dbm_multiply_opencl_print(stderr, "sg", (int)wgsize2);
dbm_multiply_opencl_print(stderr, "lu", lu);
fprintf(stderr, " ms=%.1f\n", 1E3 * ds);
} else {
fprintf(stderr, "INFO ACC/LIBDBM: DBM-kernel failed to generate\n");
@ -196,6 +235,9 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, double alpha,
#else
#error "OpenCL kernel code not found!"
#endif
if (NULL != lock_memory) {
ACC_OPENCL_ACQUIRE(lock_memory);
}
result |= c_dbcsr_acc_opencl_info_devptr_lock(&adata, NULL /*lock*/,
pack_a_data, 1 /*esize*/,
NULL /*amount*/, &iadata);
@ -206,11 +248,17 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, double alpha,
shard_c_data, 1 /*esize*/,
NULL /*amount*/, &icdata);
result |= c_dbcsr_acc_opencl_info_devptr_lock(
&batch, NULL /*lock*/, tasks /*batch*/, sizeof(dbm_task_t), &work_tasks,
&ibatch);
&batch, NULL /*lock*/, params /*batch*/, sizeof(int) * stride,
&work_tasks, &ibatch);
if (NULL != lock_memory) {
ACC_OPENCL_RELEASE(lock_memory);
}
assert(0 == iadata && 0 == ibdata && 0 == icdata);
result |= clSetKernelArg(kernel, 0, sizeof(cl_double), &alpha);
result |= clSetKernelArg(kernel, 1, sizeof(cl_int), &ibatch);
if (NULL != event || 1 == ndims) {
dbm_multiply_gpu_launch_info(&info, params_host, ntasks, param_format);
}
if (1 < ndims) { /* DBM_MULTIPLY_GEN */
const cl_uint zero = 0;
assert(0 != wgsize[1] && 0 != wgsize[1] && 0 != wgsize[2]);
@ -232,24 +280,25 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, double alpha,
work_size[0] = (0 < wgsize[0] ? LIBXSMM_UP(size, wgsize[0]) : size);
result |= clSetKernelArg(kernel, 2, sizeof(cl_int), &ntasks);
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);
result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 7, cdata.memory);
result |= clSetKernelArg(kernel, 4, sizeof(cl_int), &param_format);
result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 5, batch.memory);
result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 6, adata.memory);
result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 7, bdata.memory);
result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 8, cdata.memory);
}
result |= clEnqueueNDRangeKernel(str->queue, kernel, ndims, NULL, work_size,
0 < wgsize[0] ? wgsize : NULL,
0 /*num_wait*/, NULL /*wait_list*/, event);
}
#if defined(OPENCL_LIBSMM_PFORMAT)
#if defined(OPENCL_LIBSMM_PFORMAT) && (0 < OPENCL_LIBSMM_PFORMAT)
else { /* homogeneous */
const int pzero = 0, pbase = 3, pnext = 6;
const int param_format = pzero | (pbase << 8) | (pnext << 16);
result |= opencl_libsmm_acc_process(
NULL /*tasks_host*/, &tasks->m, ntasks, dbcsr_type_real_8, pack_a_data,
params_host, params, ntasks, dbcsr_type_real_8, pack_a_data,
pack_b_data, shard_c_data, info.max_m, info.max_n, info.max_k,
max_kernel_dim, 1 /*homogeneous*/, stream, NULL /*c_stream*/,
param_format, event);
dbm_multiply_opencl_smm, 1 /*homogeneous*/, stream, NULL /*c_stream*/,
info.max_m | info.max_n << OPENCL_LIBSMM_PFORMAT |
(info.max_k << (OPENCL_LIBSMM_PFORMAT * 2)),
event);
}
#endif
if (NULL != event && NULL != *event && EXIT_SUCCESS == result &&
@ -258,12 +307,12 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, double alpha,
const DBM_TIMER_TICKINT stop = DBM_TIMER_TICK();
const double dhost = DBM_TIMER_DIFF(start, stop);
const double diter = (0 < start2 ? DBM_TIMER_DIFF(start, start2) : dhost);
#if defined(OPENCL_LIBSMM_PFORMAT)
const char *const kind = (0 == max_kernel_dim ? "DBM" : "SMM");
#if defined(OPENCL_LIBSMM_PFORMAT) && (0 < OPENCL_LIBSMM_PFORMAT)
const char *const kind = (0 >= dbm_multiply_opencl_smm ? "DBM" : "SMM");
#else
const char *const kind = "DBM";
#endif
const int pure = (100 * (ntasks - info.mnk_count) + ntasks - 1) / ntasks;
const int pure = (100 * (ntasks - info.mnk_changes) + ntasks - 1) / ntasks;
double dkrnl = dhost, dtotl;
if (c_dbcsr_acc_opencl_timer_host == config->timer) {
cl_ulong begin = 0, end = 0;
@ -284,6 +333,18 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, double alpha,
1E+3 * diter, 1E+3 * dhost, 1E+3 * dkrnl,
1E-9 * info.max_m * info.max_n * info.max_k * ntasks / dtotl);
}
return result;
}
void dbm_multiply_gpu_launch_kernel(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) {
const int result = dbm_multiply_opencl_launch_kernel(
stream, alpha, ntasks, 0 /*param_format*/, &tasks_host->m, &tasks->m,
pack_a_data, pack_b_data, shard_c_data);
OFFLOAD_CHECK(result);
}

104
src/dbm/dbm_multiply_opencl.cl
View file

@ -12,55 +12,53 @@
#define SINT short
#define X(T, I) (T)->I /* task can be taken by value or by pointer */
#define XC(T) X(T, offset_c)
#define XK(T) (SINT) X(T, k)
#if !defined(CLINEAR)
#define XA(T) X(T, offset_a)
#define XB(T) X(T, offset_b)
#define XM(T) (SINT) X(T, m)
#define XN(T) (SINT) X(T, n)
#define XM(T) T[0]
#define XN(T) T[1]
#define XI IDT
#else
#define XA(T) X(T, offset_b)
#define XB(T) X(T, offset_a)
#define XM(T) (SINT) X(T, n)
#define XN(T) (SINT) X(T, m)
#define XM(T) T[1]
#define XN(T) T[0]
#define XI IDX
#endif
#define DBM_MULTIPLY_STORE(ALPHA, TASK, CMAT, CVEC, M, N0, N1) \
do { /* CMAT atomically accumulates CVEC */ \
#define XK(T) T[2]
#define XA(T, IBASE) (XM(T) - IBASE)
#define XB(T, IBASE) (XN(T) - IBASE)
#define XC(T, IBASE) (XK(T) - IBASE)
#define DBM_MULTIPLY_STORE(ALPHA, IBASE, SHIFT, SHAPE, C, CVEC, M, N0, N1) \
do { /* C atomically accumulates CVEC */ \
UNROLL_AUTO for (SINT n = 0; n < (N1); ++n) { /* flush to global */ \
const int idx = XI(M, n + (N0), XM(TASK), XN(TASK)) + XC(TASK); \
ACCUMULATE((CMAT) + idx, (ALPHA) * (CVEC)[n]); \
const int im = XI(M, n + (N0), XM(SHAPE), XN(SHAPE)); \
ACCUMULATE((C) + XC(SHIFT, IBASE) + im, (ALPHA) * (CVEC)[n]); \
} \
} while (0)
#define DBM_MULTIPLY_KERNEL(TASK, AMAT, BMAT, CVEC, M, N0, BN, BK) \
#define DBM_MULTIPLY_KERNEL(IBASE, SHIFT, SHAPE, A, B, CVEC, M, N0, BN, BK) \
do { /* CVEC accumulates result */ \
UNROLL(BK) for (SINT k = 0; k < XK(TASK); ++k) { \
const double a = (AMAT)[XA(TASK) + IDT(M, k, XM(TASK), XK(TASK))]; \
const int idx = IDX(k, N0, XK(TASK), XN(TASK)); \
UNROLL(BK) for (SINT k = 0; k < XK(SHAPE); ++k) { \
const int ik = IDX(k, N0, XK(SHAPE), XN(SHAPE)); \
const int ia = IDT(M, k, XM(SHAPE), XK(SHAPE)); \
const double ak = (A)[XA(SHIFT, IBASE) + ia]; \
UNROLL_AUTO for (SINT n = 0; n < (BN); ++n) { \
(CVEC)[n] = MAD(a, (BMAT)[idx + n], (CVEC)[n]); \
(CVEC)[n] = MAD(ak, (B)[ik + n], (CVEC)[n]); \
} \
} \
} while (0)
#define DBM_MULTIPLY(ALPHA, TASK, AMAT, BMAT, CMAT, CVEC, M, BN, BK) \
#define DBM_MULTIPLY(ALPHA, IBASE, SHIFT, SHAPE, A, B, C, CVEC, M, BN, BK) \
do { /* DBM_MULTIPLY_KERNEL specialized over N */ \
SINT n0 = 0, n1 = XN(TASK) - (BN); \
SINT n0 = 0, n1 = XN(SHAPE) - (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, BK); \
DBM_MULTIPLY_STORE(ALPHA, TASK, CMAT, CVEC, M, n0, BN); \
DBM_MULTIPLY_KERNEL(IBASE, SHIFT, SHAPE, A, B, CVEC, M, n0, BN, BK); \
DBM_MULTIPLY_STORE(ALPHA, IBASE, SHIFT, SHAPE, C, CVEC, M, n0, BN); \
UNROLL_FORCE(BN) for (SINT n = 0; n < (BN); ++n) { (CVEC)[n] = ZERO; } \
} \
n1 = XN(TASK) - n0; \
DBM_MULTIPLY_KERNEL(TASK, AMAT, BMAT, CVEC, M, n0, n1, BK); \
DBM_MULTIPLY_STORE(ALPHA, TASK, CMAT, CVEC, M, n0, n1); \
n1 = XN(SHAPE) - n0; \
DBM_MULTIPLY_KERNEL(IBASE, SHIFT, SHAPE, A, B, CVEC, M, n0, n1, BK); \
DBM_MULTIPLY_STORE(ALPHA, IBASE, SHIFT, SHAPE, C, CVEC, M, n0, n1); \
} while (0)
#if defined(WG) && (0 < WG)
@ -70,16 +68,14 @@ __attribute__((intel_reqd_sub_group_size(SG)))
#endif
#endif
kernel void
dbm_multiply(double alpha, int itask, int ntasks, int size,
global const dbm_task_t *tasks,
dbm_multiply(double alpha, int itask, int ntasks, int size, int param_format,
global const int *params,
#if !defined(CLINEAR)
global const double *restrict amat,
global const double *restrict bmat,
global const double *restrict a, global const double *restrict b,
#else
global const double *restrict bmat,
global const double *restrict amat,
global const double *restrict b, global const double *restrict a,
#endif
global double *restrict cmat) {
global double *restrict c) {
const int i = (int)get_global_id(0);
#if defined(SM) && (0 < SM)
local double tls[WG][BN + SM - 1], *const cvec = &tls[get_local_id(0)];
@ -90,22 +86,36 @@ dbm_multiply(double alpha, int itask, int ntasks, int size,
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];
SINT shape[3], ibase = 0, m;
int tid = i;
shape[0] = size / ntasks;
shape[1] = 0xFF & (param_format >> 8);
shape[2] = 0xFF & (param_format >> 16);
tid /= shape[0];
m = i - tid * shape[0];
if (0 == param_format) {
const int task = (itask + tid) * 6;
shape[0] = params[task + 0];
shape[1] = params[task + 1];
shape[2] = params[task + 2];
params += task + 3;
} else {
params += (itask + tid) * 3;
ibase = 1;
}
#if !defined(NDEBUG)
if (m < XM(task))
if (m < XM(shape))
#endif
{ /* valid slice (subtask) */
bmat += XB(task);
if (16 <= XK(task)) {
DBM_MULTIPLY(alpha, task, amat, bmat, cmat, cvec, m, BN, 16);
} else if (8 <= XK(task)) {
DBM_MULTIPLY(alpha, task, amat, bmat, cmat, cvec, m, BN, 8);
} else if (4 <= XK(task)) {
DBM_MULTIPLY(alpha, task, amat, bmat, cmat, cvec, m, BN, 4);
b += XB(params, ibase);
if (16 <= XK(shape)) {
DBM_MULTIPLY(alpha, ibase, params, shape, a, b, c, cvec, m, BN, 16);
} else if (8 <= XK(shape)) {
DBM_MULTIPLY(alpha, ibase, params, shape, a, b, c, cvec, m, BN, 8);
} else if (4 <= XK(shape)) {
DBM_MULTIPLY(alpha, ibase, params, shape, a, b, c, cvec, m, BN, 4);
} else {
DBM_MULTIPLY(alpha, task, amat, bmat, cmat, cvec, m, BN, 1);
DBM_MULTIPLY(alpha, ibase, params, shape, a, b, c, cvec, m, BN, 1);
}
}
}