OpenCL: support different format for parameter-stack (#901)

- Allow LIBSMM to be reused for CP2K/DBM.
- Provide opencl_libsmm_acc_process to support param_format.
- Allow host_param_stack to be a NULL-pointer.
- Signal support for OPENCL_LIBSMM_PFORMAT.
- Allow external perf-event control.
- Updated tuned parameters.

ACC_BENCH:
- Adjusted rule for MAX_KERNEL_DIM (max_kernel_dim).

Author: hfp

src/acc/acc_bench.c

@@ -54,9 +54,6 @@
 #if !defined(ELEM_TYPE)
 #  define ELEM_TYPE double
 #endif
-#if !defined(MAX_KERNEL_DIM)
-#  define MAX_KERNEL_DIM 80
-#endif
 #if !defined(ALIGNMENT)
 #  define ALIGNMENT 64
 #endif
@@ -291,6 +288,7 @@ int main(int argc, char* argv[]) {
 #else
       const int mn = m * n, mk = m * k, kn = k * n;
 #endif
+      const int max_kernel_dim = ceil(sqrt(m * n));
       int *stack_hst = NULL, *stack_dev = NULL, *trans_hst = NULL, *trans_dev = NULL;
       ELEM_TYPE *amat_hst = NULL, *bmat_hst = NULL, *cmat_hst = NULL;
       ELEM_TYPE *amat_dev = NULL, *bmat_dev = NULL, *cmat_dev = NULL;
@@ -353,7 +351,7 @@ int main(int argc, char* argv[]) {
       PRINTF(
         "%s%s%i %i %i %i %i %i %i %i\n", 0 < argc ? argv[0] : "", 0 < argc ? " " : "", nrepeat, stack_size, m, n, k, nc, na, nb);
       PRINTF("typename (id=%i): %s\n", DBCSR_TYPE(ELEM_TYPE), DBCSR_STRINGIFY(ELEM_TYPE));
-      if (MAX_KERNEL_DIM < m || MAX_KERNEL_DIM < n || MAX_KERNEL_DIM < k) {
+      if (MAX_KERNEL_DIM < max_kernel_dim) {
         fprintf(stderr, "ERROR: Matrix shape exceeds MAX_KERNEL_DIM!\n");
         result = EXIT_FAILURE;
       }
@@ -364,7 +362,7 @@ int main(int argc, char* argv[]) {
       CHECK(c_dbcsr_acc_host_mem_allocate((void**)(void*)&stack_hst, sizeof(int) * 3 * stack_size, stream), &result, check);
       CHECK(c_dbcsr_acc_host_mem_allocate((void**)(void*)&trans_hst, sizeof(int) * nb, stream), &result, check);
       CHECK(c_dbcsr_acc_stream_sync(stream), &result, check); /* ensure host-data is allocated */
-      if (NULL != amat_hst && NULL != bmat_hst && NULL != trans_hst && NULL != stack_hst) {
+      if (NULL != amat_hst && NULL != bmat_hst && NULL != trans_hst && NULL != stack_hst && EXIT_SUCCESS == result) {
         init_stack(stack_hst, stack_size, NRAND, rnd, mn, mk, kn, nc, na, nb);
 #if defined(_OPENMP)
 #  pragma omp parallel
@@ -404,7 +402,7 @@ int main(int argc, char* argv[]) {
       }
 #if defined(USE_LIBXSMM)
       CHECK(c_dbcsr_acc_stream_sync(stream), &result, check);
-      if (NULL != amat_hst && NULL != bmat_hst && NULL != stack_hst) {
+      if (NULL != amat_hst && NULL != bmat_hst && NULL != stack_hst && EXIT_SUCCESS == result) {
         const size_t size = (sizeof(ELEM_TYPE) * (mk * na + kn * nb) + sizeof(int) * 3 * stack_size) * nrepeat_h2d;
         duration = libxsmm_timer_duration(start, libxsmm_timer_tick());
         perf_h2d = size / (duration * (1ULL << 30));
@@ -414,17 +412,17 @@ int main(int argc, char* argv[]) {
 #if defined(TRANSPOSE) && defined(VALIDATE)
       /* warmup execution and prebuild transpose-kernel */
       for (r = 0; r < warmup / 2; ++r) {
-        CHECK(libsmm_acc_transpose(trans_dev, 0 /*offset*/, nb, bmat_dev, DBCSR_TYPE(ELEM_TYPE), k, n, MAX_KERNEL_DIM, stream),
+        CHECK(libsmm_acc_transpose(trans_dev, 0 /*offset*/, nb, bmat_dev, DBCSR_TYPE(ELEM_TYPE), k, n, max_kernel_dim, stream),
           &result, check);
-        CHECK(libsmm_acc_transpose(trans_dev, 0 /*offset*/, nb, bmat_dev, DBCSR_TYPE(ELEM_TYPE), n, k, MAX_KERNEL_DIM, stream),
+        CHECK(libsmm_acc_transpose(trans_dev, 0 /*offset*/, nb, bmat_dev, DBCSR_TYPE(ELEM_TYPE), n, k, max_kernel_dim, stream),
           &result, check);
       }
 #  if defined(USE_LIBXSMM)
       CHECK(c_dbcsr_acc_stream_sync(stream), &result, check);
       start = libxsmm_timer_tick();
 #  endif
       /* to perform NN-SMMs on the device, all B-matrices are transposed upfront (SMM-kernel is limited to NT) */
-      CHECK(libsmm_acc_transpose(trans_dev, 0 /*offset*/, nb, bmat_dev, DBCSR_TYPE(ELEM_TYPE), k, n, MAX_KERNEL_DIM, stream),
+      CHECK(libsmm_acc_transpose(trans_dev, 0 /*offset*/, nb, bmat_dev, DBCSR_TYPE(ELEM_TYPE), k, n, max_kernel_dim, stream),
         &result, check);
 #  if defined(USE_LIBXSMM)
       CHECK(c_dbcsr_acc_stream_sync(stream), &result, check);
@@ -434,7 +432,7 @@ int main(int argc, char* argv[]) {
       /* warmup execution and prebuild SMM-kernel */
       for (r = 0; r < warmup; ++r) {
         CHECK(libsmm_acc_process(stack_hst, stack_dev, stack_size, DBCSR_TYPE(ELEM_TYPE), amat_dev, bmat_dev, cmat_dev, m, n, k,
-                MAX_KERNEL_DIM, 1 /*homogeneous*/, stream, stream),
+                max_kernel_dim, 1 /*homogeneous*/, stream, stream),
           &result, check);
       }
       CHECK(c_dbcsr_acc_memset_zero(cmat_dev, 0 /*offset*/, sizeof(ELEM_TYPE) * mn * nc, stream), &result, check);
@@ -445,28 +443,30 @@ int main(int argc, char* argv[]) {
       for (r = 0; r < nrepeat; ++r) {
         /* GPU-kernel is limited to C += Ai * Bi^T, i.e., NT (for NN, all Bi must be transposed upfront) */
         CHECK(libsmm_acc_process(stack_hst, stack_dev, stack_size, DBCSR_TYPE(ELEM_TYPE), amat_dev, bmat_dev, cmat_dev, m, n, k,
-                MAX_KERNEL_DIM, 1 /*homogeneous*/, stream, stream),
+                max_kernel_dim, 1 /*homogeneous*/, stream, stream),
           &result, check);
       }
 #if defined(USE_LIBXSMM)
       CHECK(c_dbcsr_acc_stream_sync(stream), &result, check);
       duration = libxsmm_timer_duration(start, libxsmm_timer_tick());
-      if (0 < duration && EXIT_SUCCESS == result) {
+      if (EXIT_SUCCESS == result) {
+        if (0 < duration) {
 #  if defined(TRANSPOSE) && defined(VALIDATE)
-        PRINTF("transpose: %.2g ms %.1f GFLOPS/s\n", 1000.0 * (duration + transpose) / (nrepeat * nrepeat_smm),
-          1E-9 * ((size_t)2 * m * n * k * stack_size * nrepeat * nrepeat_smm) / (duration + transpose));
+          PRINTF("transpose: %.2g ms %.1f GFLOPS/s\n", 1000.0 * (duration + transpose) / (nrepeat * nrepeat_smm),
+            1E-9 * ((size_t)2 * m * n * k * stack_size * nrepeat * nrepeat_smm) / (duration + transpose));
 #  endif
-        perf_dev = 1E-9 * ((size_t)2 * m * n * k * stack_size * nrepeat * nrepeat_smm) / duration;
-        PRINTF("device: %.2g ms %.1f GFLOPS/s\n", 1000.0 * duration / (nrepeat * nrepeat_smm), perf_dev);
-      }
-      else {
+          perf_dev = 1E-9 * ((size_t)2 * m * n * k * stack_size * nrepeat * nrepeat_smm) / duration;
+          PRINTF("device: %.2g ms %.1f GFLOPS/s\n", 1000.0 * duration / (nrepeat * nrepeat_smm), perf_dev);
+        }
+        else {
 #  if defined(TRANSPOSE)
-        PRINTF("transpose: 0 ms 0 GFLOPS/s\n");
+          PRINTF("transpose: 0 ms 0 GFLOPS/s\n");
 #  endif
-        PRINTF("device: 0 ms 0 GFLOPS/s\n");
+          PRINTF("device: 0 ms 0 GFLOPS/s\n");
+        }
       }
 #  if defined(VALIDATE)
-      {
+      if (EXIT_SUCCESS == result) {
         ELEM_TYPE* const gold_hst = (ELEM_TYPE*)(0 != check ? libxsmm_malloc(sizeof(ELEM_TYPE) * mn * nc) : NULL);
         /* determine host's performance independent of current result code/status */
         if (NULL != gold_hst && NULL != amat_hst && NULL != bmat_hst && NULL != stack_hst) {

src/acc/libsmm_acc/libsmm_acc_benchmark.cpp

@@ -323,7 +323,7 @@ int libsmm_acc_benchmark(
 //===========================================================================
 int libsmm_acc_benchmark_transpose_(int n_stack, int* stack, int* d_stack, double* mat, double* mat_trs, double* d_mat, int n,
   int mat_m, int mat_n, ACC_DRV(event) start, ACC_DRV(event) stop, char** kernel_descr, TransposeLauncher* launcher) {
-  if (mat_m > MAX_KERNEL_DIM || mat_n > MAX_KERNEL_DIM) {
+  if ((mat_m * mat_n) > (MAX_KERNEL_DIM * MAX_KERNEL_DIM)) {
     printf("Cannot transpose matrices with dimensions above %i, got (%i x %i)\n", MAX_KERNEL_DIM, mat_m, mat_n);
     exit(1);
   }

src/acc/opencl/smm/kernels/multiply.cl

@@ -53,12 +53,6 @@
 #if !defined(SINT) /* covers matrix shape */
 #  define SINT signed char
 #endif
-
-#if defined(SLM_P) && (1 < BS)
-#  define IDXBASE 0
-#else
-#  define IDXBASE 1
-#endif
 #if !defined(REPEAT)
 #  define REPEAT 1
 #endif
@@ -76,20 +70,23 @@ __attribute__((reqd_work_group_size(WG, 1, 1)))
 __attribute__((intel_reqd_sub_group_size(SG)))
 #endif
 kernel void
-FN(global T* restrict cdata, GLOBAL const T* restrict adata, GLOBAL const T* restrict bdata,
+FN(global T* restrict cdata, GLOBAL const T* restrict adata, GLOBAL const T* restrict bdata, GLOBAL const int* restrict param_stack,
 #if (1 < BS)
-  GLOBAL const int* restrict param_stack, int stack_size, int bs) {
+  int param_format, int stack_size, int bs) {
   const int gid = get_group_id(0), idx = get_local_id(0);
 #else
-  GLOBAL const int* restrict param_stack) {
+  int param_format) {
   const int gid = get_group_id(0), idx = get_local_id(0), bs = 1;
 #endif
-  /* indexes given by param_stack are one-based (Fortran) */
-  GLOBAL const int* restrict pbase = param_stack + gid * (3 * bs);
+  const int pzero = (param_format) & 255;
+  const int pbase = (param_format >> 8) & 255;
+  const int pnext = (param_format >> 16) & 255;
+  /* param_stack/indexes can be one-based (Fortran) depending on param_format */
+  GLOBAL const int* restrict param_base = param_stack + gid * (pnext * bs);
 #if defined(SLM_P) && (1 < BS)
   local int params[3 * BS]; /* bs <= BS */
 #else
-  GLOBAL const int* restrict params = pbase;
+  GLOBAL const int* restrict params = param_base;
 #endif
 #if defined(SLM_A)
 #  if (1 != BK || BM < SM || 1 != BN)
@@ -143,9 +140,7 @@ FN(global T* restrict cdata, GLOBAL const T* restrict adata, GLOBAL const T* res
   int c0;
 #  if defined(SLM_C)
   local T cnm[SN][SM + SLM_C - 1]; /* tile in SLM */
-  for (SINT n = (SINT)idx; n < SN; n += WG) {
-    UNROLL_FORCE(SM) for (SINT m = 0; m < SM; ++m) cnm[n][m] = ZERO;
-  }
+  UNROLL_AUTO for (SINT n = (SINT)idx; n < SN; n += WG) { UNROLL_FORCE(SM) for (SINT m = 0; m < SM; ++m) cnm[n][m] = ZERO; }
 #  elif (BM < SM || 1 != BN)
 #    if (1 != BN)
   UNROLL(BN)
@@ -158,15 +153,21 @@ FN(global T* restrict cdata, GLOBAL const T* restrict adata, GLOBAL const T* res
   UNROLL_FORCE(SM) for (SINT m = 0; m < SM; ++m) cnm[m] = ZERO;
 #  endif
 #  if defined(SLM_P)
-  UNROLL_FORCE(3 * BS) for (int i = idx; i < (3 * batchsize); i += WG) params[i] = pbase[i] - 1;
+  UNROLL_AUTO for (int i = idx; i < batchsize; i += WG) {
+    UNROLL_FORCE(3) for (int j = 0; j < 3; ++j) { params[3 * i + j] = param_base[pnext * i + pbase + j] - pzero; }
+  }
 #  endif
 #  if defined(BARRIER) && (MAX(1, SG) < WG) && (defined(SLM_C) || defined(SLM_P))
   BARRIER(CLK_LOCAL_MEM_FENCE);
 #  endif
 #  if (WRK < WG)
   if (WRK <= idx) return; /* WRK <= idx */
 #  endif
-  c0 = params[2] - IDXBASE;
+#  if defined(SLM_P)
+  c0 = params[pbase + 2];
+#  else
+  c0 = params[pbase + 2] - pzero;
+#  endif
 #  if defined(BSC) && (1 != BK) && (1 != UM)
   UNROLL_OUTER(REPEAT * BS)
 #  else
@@ -179,21 +180,27 @@ FN(global T* restrict cdata, GLOBAL const T* restrict adata, GLOBAL const T* res
   for (int item = 0; item < (REPEAT * batchsize); ++item) {
     const int i = item;
 #  endif
-    const int a0 = params[3 * i] - IDXBASE, b0 = params[3 * i + 1] - IDXBASE;
-    const int c1 = ((i + 1) < batchsize ? (params[3 * i + 5] - IDXBASE) : -1);
+#  if defined(SLM_P)
+    const int idxstride = 3, idxbase = 0;
+#  else
+    const int idxstride = pnext, idxbase = pzero;
+#  endif
+    const int a0 = params[idxstride * i + pbase] - idxbase, b0 = params[idxstride * i + pbase + 1] - idxbase;
+    const int c1 = ((i + 1) < batchsize ? (params[idxstride * i + pbase + idxstride + 2] - idxbase) : -1);
 #else
 #  if (WRK < WG)
   if (WRK > idx) /* WRK > idx */
 #  endif
   {
-    const int a0 = params[0] - IDXBASE, b0 = params[1] - IDXBASE, c0 = params[2] - IDXBASE;
+    const int a0 = params[pbase + 0] - pzero, b0 = params[pbase + 1] - pzero;
+    const int c0 = params[pbase + 2] - pzero;
 #endif
 
 #if defined(SLM_A) && (1 != BK || BM < SM || 1 != BN)
     { /* copy or transpose A-matrix into SLM */
       int m = idx;
 #  if (WRK != SM)
-      for (; m < SM; m += WRK)
+      UNROLL_AUTO for (; m < SM; m += WRK)
 #  endif
       {
         UNROLL_FORCE(SK) for (SINT k = 0; k < SK; ++k) amk[m][k] = ADX(m, k);
@@ -205,7 +212,7 @@ FN(global T* restrict cdata, GLOBAL const T* restrict adata, GLOBAL const T* res
     { /* copy or transpose B-matrix into SLM */
       int n = idx;
 #  if (WRK != SN)
-      for (; n < SN; n += WRK)
+      UNROLL_AUTO for (; n < SN; n += WRK)
 #  endif
       {
         UNROLL(SK) for (SINT k = 0; k < SK; ++k) bnk[n][k] = BDX(k, n);
@@ -252,10 +259,10 @@ FN(global T* restrict cdata, GLOBAL const T* restrict adata, GLOBAL const T* res
       T cnm[BN]; /* row */
       UNROLL_FORCE(BN) for (SINT n = 0; n < BN; ++n) cnm[n] = ZERO;
 #    endif
-      UNROLL(BM)
 #    if (SM % BM)
-      for (SINT bm = 0, m = m0; bm < BM && m < SM; m = ++bm + m0)
+      UNROLL_AUTO for (SINT bm = 0, m = m0; bm < BM && m < SM; m = ++bm + m0)
 #    else
+      UNROLL(BM)
       for (SINT bm = 0, m = m0; bm < BM; m = ++bm + m0)
 #    endif
       { /* general BK, A in registers */
@@ -337,10 +344,10 @@ FN(global T* restrict cdata, GLOBAL const T* restrict adata, GLOBAL const T* res
 #    else
             const T b = BNK(n, k);
 #    endif
-            UNROLL_FORCE(BM)
 #    if (SM % BM)
-            for (SINT bm = 0, m = m0; bm < BM && m < SM; m = ++bm + m0)
+            UNROLL_AUTO for (SINT bm = 0, m = m0; bm < BM && m < SM; m = ++bm + m0)
 #    else
+            UNROLL_FORCE(BM)
             for (SINT bm = 0, m = m0; bm < BM; m = ++bm + m0)
 #    endif
             {
@@ -387,10 +394,10 @@ FN(global T* restrict cdata, GLOBAL const T* restrict adata, GLOBAL const T* res
         T cnm[BM]; /* column-block */
         UNROLL_FORCE(BM) for (SINT m = 0; m < BM; ++m) cnm[m] = ZERO;
 #    endif
-        UNROLL(BM)
 #    if (SM % BM)
-        for (SINT bm = 0, m = m0; bm < BM && m < SM; m = ++bm + m0)
+        UNROLL_AUTO for (SINT bm = 0, m = m0; bm < BM && m < SM; m = ++bm + m0)
 #    else
+        UNROLL(BM)
         for (SINT bm = 0, m = m0; bm < BM; m = ++bm + m0)
 #    endif
         {
@@ -440,7 +447,7 @@ FN(global T* restrict cdata, GLOBAL const T* restrict adata, GLOBAL const T* res
         const T b = BNK(idx, k);
 #    if defined(SLM_A)
 #      if (WRK != SM)
-        for (SINT m = (SINT)idx; m < SM; m += WRK) amk[m] = ADX(m, k);
+        UNROLL_AUTO for (SINT m = (SINT)idx; m < SM; m += WRK) amk[m] = ADX(m, k);
 #      else
         amk[idx] = ADX(idx, k);
 #      endif
@@ -483,7 +490,7 @@ FN(global T* restrict cdata, GLOBAL const T* restrict adata, GLOBAL const T* res
 #    if (1 == UM)
       UNROLL_OUTER(SM)
 #    endif
-      for (; m < (SM - UM + 1); m += UM) {
+      UNROLL_AUTO for (; m < (SM - UM + 1); m += UM) {
         u = 0;
 #    if (1 < UM)
         UNROLL(UM)
@@ -571,10 +578,10 @@ FN(global T* restrict cdata, GLOBAL const T* restrict adata, GLOBAL const T* res
         if (n < SN) /* n < SN */
 #    endif
         {
-          UNROLL_FORCE(BM)
 #    if (SM % BM)
-          for (SINT bm = 0, m = m0; bm < BM && m < SM; m = ++bm + m0)
+          UNROLL_AUTO for (SINT bm = 0, m = m0; bm < BM && m < SM; m = ++bm + m0)
 #    else
+          UNROLL_FORCE(BM)
           for (SINT bm = 0, m = m0; bm < BM; m = ++bm + m0)
 #    endif
           {
@@ -605,7 +612,7 @@ FN(global T* restrict cdata, GLOBAL const T* restrict adata, GLOBAL const T* res
     { /* atomically commit C-column to global memory */
       SINT m = 0;
 #    if defined(ATOMIC_ADD2_GLOBAL)
-      for (; m < (SM - 1); m += 2) {
+      UNROLL_AUTO for (; m < (SM - 1); m += 2) {
 #      if defined(ATOMIC_INC_NZ)
         if (ZERO != CNM(idx, m) && ZERO != CNM(idx, m + 1))
 #      endif
@@ -619,7 +626,7 @@ FN(global T* restrict cdata, GLOBAL const T* restrict adata, GLOBAL const T* res
       UNROLL(SM)
 #    endif
 #    if !defined(ATOMIC_ADD2_GLOBAL) || (SM & 1)
-      for (; m < SM; ++m) {
+      UNROLL_AUTO for (; m < SM; ++m) {
 #      if defined(ATOMIC_INC_NZ)
         if (ZERO != CNM(idx, m))
 #      endif