From b8244760835530bbb76e591c474fbfa2df263580 Mon Sep 17 00:00:00 2001 From: Hans Pabst Date: Tue, 8 Apr 2025 10:37:23 +0200 Subject: [PATCH] DBM/OpenCL: DBCSR/DBM kernel interoperability (#4126) - 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). --- src/dbm/Makefile | 6 +- src/dbm/dbm_multiply_gpu_kernel.cu | 5 +- src/dbm/dbm_multiply_gpu_kernel.h | 10 +- src/dbm/dbm_multiply_opencl.c | 217 ++++++++++++++++++----------- src/dbm/dbm_multiply_opencl.cl | 104 +++++++------- 5 files changed, 207 insertions(+), 135 deletions(-) diff --git a/src/dbm/Makefile b/src/dbm/Makefile index 1b0260994a..683fa49712 100644 --- a/src/dbm/Makefile +++ b/src/dbm/Makefile @@ -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 diff --git a/src/dbm/dbm_multiply_gpu_kernel.cu b/src/dbm/dbm_multiply_gpu_kernel.cu index 43f8d16e36..e7100059a4 100644 --- a/src/dbm/dbm_multiply_gpu_kernel.cu +++ b/src/dbm/dbm_multiply_gpu_kernel.cu @@ -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, diff --git a/src/dbm/dbm_multiply_gpu_kernel.h b/src/dbm/dbm_multiply_gpu_kernel.h index 7c74561cba..542cd031a7 100644 --- a/src/dbm/dbm_multiply_gpu_kernel.h +++ b/src/dbm/dbm_multiply_gpu_kernel.h @@ -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 } diff --git a/src/dbm/dbm_multiply_opencl.c b/src/dbm/dbm_multiply_opencl.c index 41f59bb9fe..3cecc0ee97 100644 --- a/src/dbm/dbm_multiply_opencl.c +++ b/src/dbm/dbm_multiply_opencl.c @@ -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), ¶m_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); } diff --git a/src/dbm/dbm_multiply_opencl.cl b/src/dbm/dbm_multiply_opencl.cl index b1b6efc7a0..be539d023e 100644 --- a/src/dbm/dbm_multiply_opencl.cl +++ b/src/dbm/dbm_multiply_opencl.cl @@ -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); } } }