diff --git a/src/dbm/dbm_multiply_opencl.c b/src/dbm/dbm_multiply_opencl.c index d78b5d9d70..9e5e7e25cb 100644 --- a/src/dbm/dbm_multiply_opencl.c +++ b/src/dbm/dbm_multiply_opencl.c @@ -18,50 +18,116 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, const double *pack_b_data, double *shard_c_data) { static cl_kernel kernel = NULL; - static int task_split = 0; + static int ndims = 1, split = 0; + static size_t wgsize[] = {0, 0, 0}; int result = EXIT_SUCCESS, verbosity = c_dbcsr_acc_opencl_config.verbosity; 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 int max_m = mnk_range[0][1], max_n = mnk_range[1][1]; - const size_t work_tasks = ntasks, wgsize = 0; - size_t ibatch = 0, iadata = 0, ibdata = 0, icdata = 0, work_size = 0; + const size_t max_m = mnk_range[0][1], work_tasks = ntasks; + size_t work_size[] = {1, 1, 1}, ibatch = 0, iadata = 0, ibdata = 0, + icdata = 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_range[0][0] && 0 < max_m && mnk_range[0][0] <= max_m); - assert(0 < mnk_range[1][0] && 0 < max_n && mnk_range[1][0] <= max_n); + assert(0 < mnk_range[0][0] && 0 < mnk_range[0][1] && + mnk_range[0][0] <= mnk_range[0][1]); + assert(0 < mnk_range[1][0] && 0 < mnk_range[1][1] && + mnk_range[1][0] <= mnk_range[1][1]); assert(0 < mnk_range[2][0] && 0 < mnk_range[2][1] && mnk_range[2][0] <= mnk_range[2][1]); assert(NULL != str && NULL != str->queue); assert(0 < ntasks && NULL != tasks); /* creating/calling kernel must be consistent across threads */ ACC_OPENCL_ACQUIRE(c_dbcsr_acc_opencl_config.lock_main); -#if defined(OPENCL_DBM_SOURCE_MULTIPLY_OPENCL) +#if defined(OPENCL_DBM_SOURCE_MULTIPLY) if (NULL == kernel) { /* first-time check if kernel is present */ - char params[ACC_OPENCL_BUFFERSIZE] = "-DLU=0 -DBN=4"; - const char *const flags = "-cl-fast-relaxed-math -cl-denorms-are-zero"; - const char *extensions[] = {NULL, NULL}; - const size_t nextensions = sizeof(extensions) / sizeof(*extensions); - const size_t offset = strlen(params); const libxsmm_timer_tickint start = libxsmm_timer_tick(); - const int nchar = c_dbcsr_acc_opencl_flags_atomics( + char params[ACC_OPENCL_BUFFERSIZE] = + "-cl-fast-relaxed-math -cl-denorms-are-zero"; + const char *const gen_env = getenv("DBM_MULTIPLY_GEN"); + const char *const xf_env = getenv("DBM_MULTIPLY_XF"); + const char *const lu_env = getenv("DBM_MULTIPLY_LU"); + const char *const bn_env = getenv("DBM_MULTIPLY_BN"); + const int gpu = + (CL_DEVICE_TYPE_GPU == c_dbcsr_acc_opencl_config.device.type); + const int gen = (NULL == gen_env ? 0 /*default*/ : atoi(gen_env)); + const int xf = (NULL == xf_env ? -1 /*default*/ : atoi(xf_env)); + const int lu = LIBXSMM_CLMP(NULL == lu_env ? 0 : atoi(lu_env), -2, 1); + int bn = (NULL == bn_env ? 8 : atoi(bn_env)); + const char *extensions[] = {NULL, NULL}, *flags = NULL; + size_t nextensions = sizeof(extensions) / sizeof(*extensions); + const size_t wgsize0 = c_dbcsr_acc_opencl_config.device.wgsize[0]; + const size_t wgsize1 = c_dbcsr_acc_opencl_config.device.wgsize[1]; + size_t wgsize2 = c_dbcsr_acc_opencl_config.device.wgsize[2]; + size_t offset = (0 == c_dbcsr_acc_opencl_config.debug ? strlen(params) : 0); + offset += (size_t)c_dbcsr_acc_opencl_flags_atomics( &c_dbcsr_acc_opencl_config.device, c_dbcsr_acc_opencl_atomic_fp_64, - extensions, nextensions, params + offset, sizeof(params) - offset); - if (0 < nchar && sizeof(params) > (offset + nchar)) { - const char *const task_split_env = getenv("LIBDBM_TASK_SPLIT"); - task_split = (NULL == task_split_env ? 1 : atoi(task_split_env)); - result |= c_dbcsr_acc_opencl_kernel( - 0 /*source_is_file*/, OPENCL_DBM_SOURCE_MULTIPLY_OPENCL, - "dbm_multiply", params, - 0 == c_dbcsr_acc_opencl_config.debug ? flags : NULL, NULL /*try*/, - NULL /*try_ok*/, extensions, nextensions, &kernel); - if (2 <= verbosity || 0 > verbosity) { - if (EXIT_SUCCESS == result) { - fprintf(stderr, "INFO ACC/LIBDBM: DBM-kernel ms=%.1f\n", - 1E3 * libxsmm_timer_duration(start, libxsmm_timer_tick())); - } else { - fprintf(stderr, "INFO ACC/LIBDBM: DBM-kernel failed to generate\n"); + extensions, &nextensions, params + offset, sizeof(params) - offset); + if (2 <= gen || (0 != gen && 0 != wgsize2 /*subgroups*/ && + 2 <= *c_dbcsr_acc_opencl_config.device.std_level && + NULL != extensions[1] && + NULL != strstr(extensions[1], "cl_ext_float_atomics"))) { + offset += + (size_t)LIBXSMM_SNPRINTF(params + offset, sizeof(params) - offset, + " -DDBM_MULTIPLY_OPENCL_GEN"); + if (0 != c_dbcsr_acc_opencl_config.device.intel && 0 != xf) { + flags = "-cl-intel-256-GRF-per-thread"; + } + wgsize[1] = wgsize[2] = 1; + wgsize[0] = 16; + ndims = 3; + } else { + const char *const split_env = getenv("DBM_MULTIPLY_SPLIT"); + const char *const wg_env = getenv("DBM_MULTIPLY_WG"); + split = (NULL == split_env ? 1 /*default*/ : atoi(split_env)); + wgsize[0] = + (NULL == wg_env ? (1 != split ? (wgsize1 * LIBXSMM_ABS(split)) : 0) + : strtoul(wg_env, NULL, 10)); + if (0 != split && 1 != split && (bn * bn) > (int)wgsize[0]) { + wgsize[0] = bn * bn; + } + if (0 != split && 0 != wgsize2 && 0 < wgsize[0]) { /* subgroups */ + if (LIBXSMM_DELTA(wgsize[0], wgsize1) <= + LIBXSMM_DELTA(wgsize[0], wgsize2)) { /* select SG-size */ + wgsize2 = wgsize1; } + wgsize[0] = LIBXSMM_UP(wgsize[0], wgsize2); + } else { + wgsize[0] = LIBXSMM_UP(wgsize[0], wgsize1); + wgsize2 = 0; + } + wgsize[0] = LIBXSMM_CLMP(wgsize[0], 0, wgsize0); + if (NULL == bn_env && 0 != split && 1 != split && + (bn * bn) < (int)wgsize[0]) { + bn = libxsmm_isqrt2_u32(wgsize[0]); + } + bn = LIBXSMM_CLMP(bn, 4, 32); + offset += (size_t)LIBXSMM_SNPRINTF( + params + offset, sizeof(params) - offset, + " %s -DSPLIT=%i -DBN=%i -DWG=%i -DSG=%i -DLU=%i", + 0 != gpu ? "-DGPU" : "", split, bn, (int)wgsize[0], (int)wgsize2, lu); + if (0 != c_dbcsr_acc_opencl_config.device.intel && 0 < xf) { + flags = "-cl-intel-256-GRF-per-thread"; + } + } + result |= (sizeof(params) > 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, nextensions, + &kernel); + if (2 <= verbosity || 0 > verbosity) { + if (EXIT_SUCCESS == result) { + const double d = libxsmm_timer_duration(start, libxsmm_timer_tick()); + fprintf(stderr, "INFO ACC/LIBDBM: DBM-kernel gpu=%i", gpu); + if (0 == gen) { + fprintf(stderr, " split=%i lu=%i bn=%i", split, lu, bn); + } else { /* generated kernel */ + fprintf(stderr, " gen=%i", gen); + } + fprintf(stderr, " wg=%i sg=%i ms=%.1f\n", (int)wgsize[0], (int)wgsize2, + 1E3 * d); + } else { + fprintf(stderr, "INFO ACC/LIBDBM: DBM-kernel failed to generate\n"); } } } @@ -81,18 +147,47 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, &batch, NULL /*lock*/, tasks /*batch*/, sizeof(dbm_task_t), &work_tasks, &ibatch); assert(0 == iadata && 0 == ibdata && 0 == icdata); - work_size = (0 != task_split ? (work_tasks * max_m) : work_tasks); result |= clSetKernelArg(kernel, 0, sizeof(cl_double), &alpha); result |= clSetKernelArg(kernel, 1, sizeof(cl_int), &ibatch); - result |= clSetKernelArg(kernel, 2, sizeof(cl_int), &ntasks); - result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 3, batch.memory); - result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 4, adata.memory); - result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 5, bdata.memory); - result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 6, cdata.memory); - result |= clEnqueueNDRangeKernel(str->queue, kernel, 1 /*work_dim*/, - NULL /*offset*/, &work_size, - 0 != wgsize ? &wgsize : NULL, 0 /*num_wait*/, - NULL /*wait_list*/, perf_event); + if (1 < ndims) { /* generated kernel */ + const cl_uint zero = 0; + assert(0 != wgsize[1] && 0 != wgsize[1] && 0 != wgsize[2]); + work_size[0] = 16; + assert(1 == work_size[1]); + work_size[2] = work_tasks; + result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 2, batch.memory); + result |= clSetKernelArg(kernel, 3, sizeof(cl_uint), &zero /*shape*/); + result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 4, adata.memory); + result |= clSetKernelArg(kernel, 5, sizeof(cl_uint), &zero /*A_shape0*/); + result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 6, bdata.memory); + result |= clSetKernelArg(kernel, 7, sizeof(cl_uint), &zero /*B_shape0*/); + result |= c_dbcsr_acc_opencl_set_kernel_ptr(kernel, 8, cdata.memory); + result |= clSetKernelArg(kernel, 9, sizeof(cl_uint), &zero /*C_shape0*/); + } else { + result |= clSetKernelArg(kernel, 2, sizeof(cl_int), &ntasks); + if (0 != split) { + if (1 == split || 0 == wgsize[0]) { + 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]); + } + } else { + work_size[0] = work_tasks * wgsize[0]; + result |= clSetKernelArg(kernel, 3, sizeof(cl_int), work_size); + } + } else { + work_size[0] = work_tasks; + result |= clSetKernelArg(kernel, 3, sizeof(cl_int), work_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 |= clEnqueueNDRangeKernel( + str->queue, kernel, ndims, NULL, work_size, 0 < wgsize[0] ? wgsize : NULL, + 0 /*num_wait*/, NULL /*wait_list*/, perf_event); if (NULL != perf_event && EXIT_SUCCESS == result) { cl_ulong begin = 0, end = 0; clWaitForEvents(1, perf_event); @@ -101,13 +196,14 @@ void dbm_multiply_gpu_launch_kernel(const offloadStream_t stream, result |= clGetEventProfilingInfo(*perf_event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL); if (EXIT_SUCCESS == result) { - const double duration_ms = 1E-6 * LIBXSMM_DELTA(begin, end); + const double duration_ns = LIBXSMM_DELTA(begin, end); const double gflops = - (1E-6 * max_m * max_n * mnk_range[2][1] * ntasks) / duration_ms; + (max_m * mnk_range[1][1] * mnk_range[2][1] * ntasks) / duration_ns; fprintf(stderr, "INFO ACC/LIBDBM: DBM-kernel mnk=%ix%ix%i " "ntasks=%i gflops=%.1f ms=%.2g\n", - max_m, max_n, mnk_range[2][1], ntasks, gflops, duration_ms); + mnk_range[0][1], mnk_range[1][1], mnk_range[2][1], ntasks, gflops, + 1E-6 * duration_ns); } } ACC_OPENCL_RELEASE(c_dbcsr_acc_opencl_config.lock_main); diff --git a/src/dbm/dbm_multiply_opencl.cl b/src/dbm/dbm_multiply_opencl.cl index 3c97883a40..c192055725 100644 --- a/src/dbm/dbm_multiply_opencl.cl +++ b/src/dbm/dbm_multiply_opencl.cl @@ -4,74 +4,209 @@ /* */ /* SPDX-License-Identifier: BSD-3-Clause */ /*----------------------------------------------------------------------------*/ +#if defined(DBM_MULTIPLY_OPENCL_GEN) +#include "dbm_multiply_opencl.ir.h" +#else #include "../../exts/dbcsr/src/acc/opencl/common/opencl_atomics.h" #include "dbm_multiply_internal.h" -#define IDX(I, J, K, M, N) ((I) * (N) + (J) + (K)) -#define IDT(I, J, K, M, N) IDX(J, I, K, N, M) +#if defined(GPU) && defined(WG) && (0 < WG) && (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 DBM_MULTIPLY_KERNEL(TASK, A, B, CV, M, N0, N1, UNROLL_N, UNROLL_K) \ - do { \ - UNROLL_K \ - for (int k = 0; k < (TASK).k; ++k) { \ - const int ia = IDT(M, k, (TASK).offset_a, (TASK).m, (TASK).k); \ - const double a = (A)[ia]; \ - UNROLL_N \ - for (int n = 0; n < (N1); ++n) { \ - const int ib = IDX(k, n + (N0), (TASK).offset_b, (TASK).k, (TASK).n); \ - (CV)[n] = MAD(a, (B)[ib], (CV)[n]); \ +#define SINT short + +#define DIVUP(A, B) (((A) + (B)-1) / (B)) +#define NUP(N, UP) (DIVUP(N, UP) * (UP)) +#define BLR(N, BN) (NUP(N, BN) - (N)) + +#define IDX(I, J, M, N) ((int)(I) * (N) + (J)) +#define IDT(I, J, M, N) IDX(J, I, N, M) +#define X(T, I) (T)->I /* task can be taken by value or by pointer */ +#define XA(T) X(T, offset_a) +#define XB(T) X(T, offset_b) +#define XC(T) X(T, offset_c) +#define XM(T) (SINT) X(T, m) +#define XN(T) (SINT) X(T, n) +#define XK(T) (SINT) X(T, k) + +#define DBM_MULTIPLY_SHM(ALPHA, TASK, AMAT, BMAT, CMAT, SHM, WG, BM, BN) \ + do { /* matrix multiplication per work-group using shared memory */ \ + local double *restrict const ashm = (SHM); \ + local double *restrict const bshm = (SHM) + (WG); \ + const int mk = XM(TASK) * XK(TASK), kn = XK(TASK) * XN(TASK); \ + const SINT tid = (SINT)get_local_id(0); \ + /* y/s can exceed BN/BM (up to BK), and x/t is fast index (up to BM/BN) */ \ + const SINT y = tid / (BM), x = tid - y * (BM), bk = (WG) / MAX(BM, BN); \ + const SINT s = tid / (BN), t = tid - s * (BN); \ + for (SINT m0 = 0; m0 < XM(TASK); m0 += (BM)) { \ + for (SINT n0 = 0; n0 < XN(TASK); n0 += (BN)) { \ + double r = ZERO; \ + UNROLL_AUTO for (SINT k0 = 0; k0 < XK(TASK); k0 += bk) { \ + if (x < (BM) && y < bk) { /* load A-tile */ \ + const int idx = IDT(m0 + x, k0 + y, XM(TASK), XK(TASK)); \ + ashm[y * (BM) + x] = (idx < mk ? (AMAT)[XA(TASK) + idx] : ZERO); \ + } \ + if (s < bk && t < (BN)) { /* load B-tile */ \ + const int idx = IDX(k0 + s, n0 + t, XK(TASK), XN(TASK)); \ + bshm[s * (BN) + t] = (idx < kn ? (BMAT)[XB(TASK) + idx] : ZERO); \ + } \ + BARRIER(CLK_LOCAL_MEM_FENCE); \ + if (x < (BM) && y < (BN)) { /* multiply tiles */ \ + UNROLL_AUTO for (SINT z = 0; z < bk; ++z) { \ + r = MAD(ashm[z * (BM) + x], bshm[z * (BN) + y], r); \ + } \ + } \ + BARRIER(CLK_LOCAL_MEM_FENCE); \ + } \ + if (x < (BM) && y < (BN)) { /* flush to global */ \ + const SINT m = m0 + x, n = n0 + y; \ + if (m < XM(TASK) && n < XN(TASK)) { \ + const int idx = IDT(m, n, XM(TASK), XN(TASK)); \ + ACCUMULATE((CMAT) + XC(TASK) + idx, (ALPHA)*r); \ + } \ + } \ } \ } \ } while (0) -#define DBM_MULTIPLY_ACCUMULATE(ALPHA, TASK, CV, C, M, BN, N0) \ - do { /* flush private accumulator to global memory using atomics */ \ - UNROLL_FORCE(BN) \ - for (int n = 0; n < (BN); ++n) { \ - const int ic = IDT(M, n + (N0), (TASK).offset_c, (TASK).m, (TASK).n); \ - ACCUMULATE((C) + ic, (ALPHA) * (CV)[n]); \ - cv[n] = ZERO; /* reset */ \ +#define DBM_MULTIPLY_KERNEL(ALPHA, TASK, AMAT, BMAT, CMAT, CVEC, M, N0, N1, K, \ + BCST) \ + UNROLL_AUTO for (SINT k = 0; k < (K); ++k) { \ + const double a = (AMAT)[XA(TASK) + IDT(M, k, XM(TASK), K)]; \ + UNROLL_AUTO for (SINT n = 0; n < (N1); ++n) { \ + const double b = (BMAT)[XB(TASK) + IDX(k, n + (N0), K, XN(TASK))]; \ + (CVEC)[n] = MAD(a, BCST(b), (CVEC)[n]); \ + } \ + } \ + 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 */ \ + } + +#define DBM_MULTIPLY(ALPHA, TASK, AMAT, BMAT, CMAT, M, BN, BCST) \ + do { /* DBM_MULTIPLY_KERNEL unrolled/specialized over N and K */ \ + double cvec[BN]; \ + SINT n0 = 0; \ + UNROLL_AUTO for (SINT n = 0; n < (BN); ++n) { cvec[n] = ZERO; } \ + if ((BN) <= XN(TASK)) { \ + if (1 < XK(TASK)) { \ + UNROLL_OUTER(1) for (; (n0 + (BN)) <= XN(TASK); n0 += (BN)) { \ + DBM_MULTIPLY_KERNEL(ALPHA, TASK, AMAT, BMAT, CMAT, cvec, M, n0, BN, \ + XK(TASK), BCST); \ + } \ + } else { /* K = 1 */ \ + UNROLL_OUTER(1) for (; (n0 + (BN)) <= XN(TASK); n0 += (BN)) { \ + DBM_MULTIPLY_KERNEL(ALPHA, TASK, AMAT, BMAT, CMAT, cvec, M, n0, BN, \ + 1, BCST); \ + } \ + } \ + } else if (1 != XK(TASK)) { /* N < BN */ \ + DBM_MULTIPLY_KERNEL(ALPHA, TASK, AMAT, BMAT, CMAT, cvec, M, 0, 1, \ + XK(TASK), BCST); \ + n0 = 1; \ + } else { /* N < BN, K = 1 */ \ + DBM_MULTIPLY_KERNEL(ALPHA, TASK, AMAT, BMAT, CMAT, cvec, M, 0, 1, 1, \ + BCST); \ + n0 = 1; \ + } \ + /*if (n0 < XN(TASK))*/ { /* handle remainder */ \ + DBM_MULTIPLY_KERNEL(ALPHA, TASK, AMAT, BMAT, CMAT, cvec, M, n0, \ + XN(TASK) - n0, XK(TASK), BCST); \ } \ } while (0) -kernel void dbm_multiply(double alpha, int itask, int ntasks, - global const dbm_task_t *tasks, - global const double *restrict am, - global const double *restrict bm, - global double *restrict cm) { - double cv[BN] = {0}; /* private accumulator */ - const int size = (int)get_global_size(0), i = (int)get_global_id(0); - - if (size != ntasks) { /* LIBDBM_TASK_SPLIT */ - const int max_m = size / ntasks, tid = i / max_m; - const dbm_task_t task = tasks[itask + min(tid, ntasks - 1)]; /* copy */ - const int m = i - tid * max_m; - if (m < task.m) { - if ((BN) < task.n) { - UNROLL_AUTO - for (int n0 = 0; n0 < task.n; n0 += (BN)) { - const int n1 = min(BN, task.n - n0); - DBM_MULTIPLY_KERNEL(task, am, bm, cv, m, n0, n1, UNROLL_FORCE(BN), - UNROLL_AUTO); - DBM_MULTIPLY_ACCUMULATE(alpha, task, cv, cm, m, BN, n0); - } - } else { /* small */ - DBM_MULTIPLY_KERNEL(task, am, bm, cv, m, 0, task.n, UNROLL_FORCE(BN), - UNROLL_FORCE(BN)); - DBM_MULTIPLY_ACCUMULATE(alpha, task, cv, cm, m, BN, 0); - } +#if defined(WG) && (0 < WG) +__attribute__((reqd_work_group_size(WG, 1, 1))) +#if defined(SG) && (0 < SG) +__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, global const double *restrict amat, + global const double *restrict bmat, global double *restrict cmat) { +#if defined(SPLIT) && (1 < SPLIT) && defined(WG) && (0 < WG) + local double shm[WG * 2]; + global const dbm_task_t *const task = &tasks[itask + get_group_id(0)]; + const SINT rmin = MIN(XM(task), XN(task)), rmax = MAX(XM(task), XN(task)); + if ((rmax - rmin) <= BN) { + if ((rmin * 4) < BN) { + DBM_MULTIPLY_SHM(alpha, task, amat, bmat, cmat, shm, WG, BN / 4, BN / 4); + } else if ((rmin * 2) < BN) { + DBM_MULTIPLY_SHM(alpha, task, amat, bmat, cmat, shm, WG, BN / 2, BN / 2); + } else { + DBM_MULTIPLY_SHM(alpha, task, amat, bmat, cmat, shm, WG, BN, BN); } - } else { /* full matrix multiplication */ - const dbm_task_t task = tasks[itask + i]; - UNROLL_OUTER(1) - for (int m = 0; m < task.m; ++m) { - UNROLL(1) - for (int n0 = 0; n0 < task.n; n0 += (BN)) { - const int n1 = min(BN, task.n - n0); - DBM_MULTIPLY_KERNEL(task, am, bm, cv, m, n0, n1, UNROLL_AUTO, - UNROLL_AUTO); - DBM_MULTIPLY_ACCUMULATE(alpha, task, cv, cm, m, BN, n0); + } else if (XM(task) <= XN(task)) { + const SINT r1 = BLR(XM(task), BN); + const SINT r2 = BLR(XM(task), BN / 2) * 2; + const SINT r3 = BLR(XM(task), BN / 4) * 4; + if (r1 <= r2) { + if (r1 <= r3) { + DBM_MULTIPLY_SHM(alpha, task, amat, bmat, cmat, shm, WG, BN, BN); + } else { + DBM_MULTIPLY_SHM(alpha, task, amat, bmat, cmat, shm, WG, BN / 4, + BN * 4); + } + } else if (r2 <= r3) { + DBM_MULTIPLY_SHM(alpha, task, amat, bmat, cmat, shm, WG, BN / 2, BN * 2); + } else { + DBM_MULTIPLY_SHM(alpha, task, amat, bmat, cmat, shm, WG, BN / 4, BN * 4); + } + } else { + const SINT r1 = BLR(XN(task), BN); + const SINT r2 = BLR(XN(task), BN / 2) * 2; + const SINT r3 = BLR(XN(task), BN / 4) * 4; + if (r1 <= r2) { + if (r1 <= r3) { + DBM_MULTIPLY_SHM(alpha, task, amat, bmat, cmat, shm, WG, BN, BN); + } else { + DBM_MULTIPLY_SHM(alpha, task, amat, bmat, cmat, shm, WG, BN * 4, + BN / 4); + } + } else if (r2 <= r3) { + DBM_MULTIPLY_SHM(alpha, task, amat, bmat, cmat, shm, WG, BN * 2, BN / 2); + } else { + DBM_MULTIPLY_SHM(alpha, task, amat, bmat, cmat, shm, WG, BN * 4, BN / 4); + } + } +#elif defined(SPLIT) && (0 != SPLIT) + const int i = (int)get_global_id(0); +#if defined(BCST_WG) + if (i < size) +#endif + { /* DBM_MULTIPLY_SPLIT */ + 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 (m < XM(task)) { /* valid task */ +#if defined(BCST_WG) + if (XM(task) <= XN(task)) { /* BCST_WG to broadcast B-values */ + DBM_MULTIPLY(alpha, task, amat, bmat, cmat, m, BN, BCST_WG); + } else +#endif + { + DBM_MULTIPLY(alpha, task, amat, bmat, cmat, m, BN, BCST_NO); } } } +#else +#if defined(BCST_WG) + if (get_global_id(0) < size) +#endif + { /* full matrix multiplication per work-item (thread) */ + global const dbm_task_t *const task = &tasks[itask + get_global_id(0)]; + UNROLL_OUTER(1) for (SINT m = 0; m < XM(task); ++m) { + DBM_MULTIPLY(alpha, task, amat, bmat, cmat, m, BN, BCST_NO); + } + } +#endif } +#endif diff --git a/tools/precommit/check_file_properties.py b/tools/precommit/check_file_properties.py index 2334082bba..cc31f9cf38 100755 --- a/tools/precommit/check_file_properties.py +++ b/tools/precommit/check_file_properties.py @@ -23,7 +23,7 @@ FLAG_EXCEPTIONS = ( r"CUDA_VERSION", r"DBM_LIBXSMM_PREFETCH", r"DBM_VALIDATE_AGAINST_DBCSR", - r"OPENCL_DBM_SOURCE_MULTIPLY_OPENCL", + r"OPENCL_DBM_SOURCE_MULTIPLY", r"FD_DEBUG", r"GRID_DO_COLLOCATE", r"INTEL_MKL_VERSION", @@ -212,7 +212,8 @@ def check_file(path: pathlib.Path) -> List[str]: if line.split()[0] not in ("#if", "#ifdef", "#ifndef", "#elif"): continue - line = line.split("//", 1)[0] + line = line.split("/*", 1)[0] # C comment + line = line.split("//", 1)[0] # C++ comment line_continuation = line.rstrip().endswith("\\") line = OP_RE.sub(" ", line) line = line.replace("defined", " ") @@ -228,6 +229,8 @@ def check_file(path: pathlib.Path) -> List[str]: flags = {flag for flag in flags if not FLAG_EXCEPTIONS_RE.match(flag)} for flag in sorted(flags): + if fn_ext == ".cl": # usually compiled at RT (no direct user-control) + continue if flag == "_OMP_H" and fn_ext == ".cu": continue if flag not in get_install_txt():