comp-chem/cp2k
1
0
Fork 0
mirror of https://github.com/cp2k/cp2k.git synced 2026-07-28 22:25:32 -04:00
Code Issues Projects Releases Packages Wiki Activity Actions

DBM/OCL: experimental single precision compute

Browse source 
- Internal option to trigger single-precision compute.
- Fixed/consider 1<sgsize as support for subgroups.
- Note: SP-compute still requires DP-capability.
This commit is contained in:
Hans Pabst 2025-08-13 09:45:56 +02:00
parent d69845916c
commit d3507fb5dc
2 changed files with 52 additions and 32 deletions
Download patch file Download diff file Expand all files Collapse all files

64
src/dbm/dbm_multiply_opencl.c
View file

@ -140,6 +140,7 @@ int dbm_multiply_opencl_launch_kernel(void *stream, double alpha, int ntasks,
"-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");
const char *const fp_env = getenv("DBM_MULTIPLY_FP");
const char *const bn_env = getenv("DBM_MULTIPLY_BN");
const char *const sm_env = getenv("DBM_MULTIPLY_SM");
const char *const wg_env = getenv("DBM_MULTIPLY_WG");
@ -150,24 +151,25 @@ int dbm_multiply_opencl_launch_kernel(void *stream, double alpha, int ntasks,
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 && 0 == param_format)
int gen = ((NULL == fp_env && NULL == bn_env && NULL == sm_env &&
NULL == wg_env && NULL == lu_env && NULL == lin_env &&
0 == param_format)
? (NULL == gen_env ? 1 /*default*/ : atoi(gen_env))
: 0);
const int precision = (NULL == fp_env ? 0 /*default*/ : atoi(fp_env));
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}, *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 sgsize = devinfo->wgsize[2];
size_t offset =
((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,
flags + offset, sizeof(flags) - offset);
if (2 <= gen ||
(0 != gen && 0 != wgsize2 /*subgroups*/ &&
2 <= *devinfo->std_level && NULL != extensions[1] &&
(0 != gen && 1 < sgsize /*subgroups*/ && 2 <= *devinfo->std_level &&
NULL != extensions[1] &&
NULL != strstr(extensions[1], "cl_ext_float_atomics"))) {
offset +=
(size_t)LIBXSMM_SNPRINTF(flags + offset, sizeof(flags) - offset,
@ -179,17 +181,17 @@ int dbm_multiply_opencl_launch_kernel(void *stream, double alpha, int ntasks,
} else {
wgsize[0] = (NULL == wg_env ? (unsigned long int)LIBXSMM_ABS(sm)
: strtoul(wg_env, NULL, 10));
if (0 != wgsize2 && 0 < wgsize[0]) { /* subgroups */
if (LIBXSMM_DELTA(wgsize[0], wgsize1) <=
LIBXSMM_DELTA(wgsize[0], wgsize2)) { /* select SG-size */
wgsize2 = wgsize1;
if (1 < sgsize && 0 < wgsize[0]) { /* subgroups */
if (LIBXSMM_DELTA(wgsize[0], devinfo->wgsize[1]) <=
LIBXSMM_DELTA(wgsize[0], sgsize)) { /* select SG-size */
sgsize = devinfo->wgsize[1];
}
wgsize[0] = LIBXSMM_UP(wgsize[0], wgsize2);
wgsize[0] = LIBXSMM_UP(wgsize[0], sgsize);
} else {
wgsize[0] = LIBXSMM_UP(wgsize[0], wgsize1);
wgsize2 = 0;
wgsize[0] = LIBXSMM_UP(wgsize[0], devinfo->wgsize[1]);
sgsize = 0;
}
wgsize[0] = LIBXSMM_CLMP(wgsize[0], 0, wgsize0);
wgsize[0] = LIBXSMM_CLMP(wgsize[0], 0, devinfo->wgsize[0]);
sm = ((0 != sm && 0 != wgsize[0])
? (LIBXSMM_ISPOT(bn * sizeof(double)) + 1)
: 0);
@ -198,31 +200,41 @@ int dbm_multiply_opencl_launch_kernel(void *stream, double alpha, int ntasks,
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);
lu, (int)wgsize[0], (int)sgsize);
if (0 != precision) {
offset +=
(size_t)LIBXSMM_SNPRINTF(flags + offset, sizeof(flags) - offset,
" -DPRECISION=%i", precision);
}
gen = 0;
}
if (0 != devinfo->intel && 0 < xf) {
options = "-cl-intel-256-GRF-per-thread";
}
result |= (sizeof(flags) > offset ? EXIT_SUCCESS : EXIT_FAILURE);
if (2 <= verbosity || 0 > verbosity || EXIT_SUCCESS != result) {
const char *const kind = (EXIT_SUCCESS == result ? "INFO" : "ERROR");
fprintf(stderr, "%s ACC/LIBDBM: DBM-kernel gpu=%i", kind, gpu);
dbm_multiply_opencl_print(stderr, "gen", gen); /* generated */
dbm_multiply_opencl_print(stderr, "lin", clinear);
dbm_multiply_opencl_print(stderr, "fp", precision);
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)sgsize);
dbm_multiply_opencl_print(stderr, "lu", lu);
fprintf(stderr, " -> ");
}
result |= c_dbcsr_acc_opencl_kernel(
0 /*source_is_file*/, OPENCL_DBM_SOURCE_MULTIPLY, "dbm_multiply",
flags, options, NULL /*try*/, NULL /*try_ok*/, extensions,
nextensions, &kernel_global);
if (2 <= verbosity || 0 > verbosity) {
if (2 <= verbosity || 0 > verbosity || EXIT_SUCCESS != result) {
if (EXIT_SUCCESS == result) {
const double ds = DBM_TIMER_DIFF(start, DBM_TIMER_TICK());
fprintf(stderr, "INFO ACC/LIBDBM: DBM-kernel gpu=%i", gpu);
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);
fprintf(stderr, "%.1f ms\n", 1E3 * ds);
} else {
fprintf(stderr, "INFO ACC/LIBDBM: DBM-kernel failed to generate\n");
fprintf(stderr, "FAILED!\n");
}
}
}

20
src/dbm/dbm_multiply_opencl.cl
View file

@ -12,6 +12,14 @@
#define SINT short
#if defined(PRECISION) && (1 == PRECISION)
#define CVT(A) convert_float(A)
#define TYPE float
#else
#define TYPE double
#define CVT(A) A
#endif
#if !defined(CLINEAR)
#define XM(T) T[0]
#define XN(T) T[1]
@ -40,9 +48,9 @@
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]; \
const TYPE ak = CVT((A)[XA(SHIFT, IBASE) + ia]); \
UNROLL_AUTO for (SINT n = 0; n < (BN); ++n) { \
(CVEC)[n] = MAD(ak, (B)[ik + n], (CVEC)[n]); \
(CVEC)[n] = MAD(ak, CVT((B)[ik + n]), (CVEC)[n]); \
} \
} \
} while (0)
@ -50,11 +58,11 @@
#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(SHAPE) - (BN); \
UNROLL_FORCE(BN) for (SINT n = 0; n < (BN); ++n) { (CVEC)[n] = ZERO; } \
UNROLL_FORCE(BN) for (SINT n = 0; n < (BN); ++n) { (CVEC)[n] = 0; } \
UNROLL_OUTER(1) for (; n0 <= n1; 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; } \
UNROLL_FORCE(BN) for (SINT n = 0; n < (BN); ++n) { (CVEC)[n] = 0; } \
} \
n1 = XN(SHAPE) - n0; \
DBM_MULTIPLY_KERNEL(IBASE, SHIFT, SHAPE, A, B, CVEC, M, n0, n1, BK); \
@ -78,9 +86,9 @@ dbm_multiply(double alpha, int itask, int ntasks, int size, int param_format,
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)];
local TYPE tls[WG][BN + SM - 1], *const cvec = &tls[get_local_id(0)];
#else
double cvec[BN];
TYPE cvec[BN];
#endif
#if defined(WG) && (0 < WG)
if (i < size)