mirror of
https://github.com/cp2k/cp2k.git
synced 2026-07-29 06:35:28 -04:00
DBM: OpenCL implementation (#3375)
* There are three main implementations 1. DBM_MULTIPLY_SPLIT=1: default kernel with smallest LOC-number uses private/register-backed buffer. This implementation is simple and performs best even on different vendor's GPUs. There is moderate unrolling/code-bloat leveraging macros. Though, there is no "old school GPU" like copy-into SLM; if there is re-use, it leverage existing caches between registers and global memory. 2. DBM_MULTIPLY_SPLIT=2..N with DBM_MULTIPLY_SPLIT=8 matching the CUDA implementation using shared/local memory. This implementation matches the existing CUDA implementation and perform equal and reasonable across GPUs. There is moderate unrolling/code-bloat leveraging macros; heuristics are revised/different from CUDA implementation (less cases). 3. DBM_MULTIPLY_GEN=1: uses https://github.com/intel/tiny-tensor-compiler and IR describing DBM's input format. The IR-file (https://github.com/hfp/cp2k/blob/master/src/dbm/dbm_multiply_opencl.ir) is not part of PR. The generated kernel-code is not yet competitive. * The amount of code in the OpenCL-kernel and and C based glue-code can be *significantly* reduced (fraction of LOC) when removing for instance implementation 2 and 3 (above mentioned). Reducing the number of LOC can be part of a future cleanup (and once code is settled). * check_file_properties.py: support C-comments, and omit checking OpenCL source code (it is usually compiled at runtime, i.e., machine-driven rather than user-controlled).
This commit is contained in:
Hans Pabst
GitHub
parent
17d3be3016
commit
25edd69807
3 changed files with 332 additions and 98 deletions
|
|
@ -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);
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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():
|
||||
|
|
|
|||
Loading…
Add table
Add a link
Reference in a new issue