Introduced ACC_OPENCL_CMEM

- Moved CMEM support from SMM to backend.
- Renamed GLOBAL to CONSTANT (kernel).

Author: hfp

src/acc/opencl/acc_opencl.c

@@ -150,6 +150,8 @@ int c_dbcsr_acc_opencl_order_devices(const void* dev_a, const void* dev_b) {
 }
 
 
+/** Setup to run prior to touching OpenCL runtime. */
+void c_dbcsr_acc_opencl_configure(void);
 void c_dbcsr_acc_opencl_configure(void) {
   const char* const env_rank = (NULL != getenv("PMI_RANK") ? getenv("PMI_RANK") : getenv("OMPI_COMM_WORLD_LOCAL_RANK"));
   const char* const env_nranks = getenv("MPI_LOCALNRANKS"); /* TODO */
@@ -762,6 +764,15 @@ int c_dbcsr_acc_finalize(void) {
 }
 
 
+int c_dbcsr_acc_opencl_use_cmem(const c_dbcsr_acc_opencl_device_t* devinfo) {
+#  if defined(ACC_OPENCL_CMEM)
+  return (0 != devinfo->size_maxalloc && devinfo->size_maxalloc <= devinfo->size_maxcmem) ? EXIT_SUCCESS : EXIT_FAILURE;
+#  else
+  return EXIT_FAILURE;
+#  endif
+}
+
+
 void c_dbcsr_acc_clear_errors(void) {}
 
 
@@ -1111,6 +1122,16 @@ int c_dbcsr_acc_opencl_set_active_device(ACC_OPENCL_LOCKTYPE* lock, int device_i
           {
             devinfo->unified = CL_FALSE;
           }
+          if (EXIT_SUCCESS !=
+              clGetDeviceInfo(active_id, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(cl_ulong), &devinfo->size_maxalloc, NULL))
+          {
+            devinfo->size_maxalloc = 0;
+          }
+          if (EXIT_SUCCESS !=
+              clGetDeviceInfo(active_id, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, sizeof(cl_ulong), &devinfo->size_maxcmem, NULL))
+          {
+            devinfo->size_maxcmem = 0;
+          }
           if (EXIT_SUCCESS != clGetDeviceInfo(active_id, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), devinfo->wgsize, NULL)) {
             devinfo->wgsize[0] = 1;
           }

src/acc/opencl/acc_opencl.h

@@ -93,6 +93,9 @@
 #if !defined(ACC_OPENCL_DELIMS)
 #  define ACC_OPENCL_DELIMS ",;"
 #endif
+#if !defined(ACC_OPENCL_CMEM) && 1
+#  define ACC_OPENCL_CMEM
+#endif
 #if !defined(ACC_OPENCL_ASYNC) && 1
 #  define ACC_OPENCL_ASYNC getenv("ACC_OPENCL_ASYNC")
 #endif
@@ -262,6 +265,8 @@ typedef struct c_dbcsr_acc_opencl_device_t {
    * smaller if an alternative SG-size exists (SG is zero if no support).
    */
   size_t wgsize[3];
+  /** Maximum size of memory allocations and constant buffer. */
+  cl_ulong size_maxalloc, size_maxcmem;
   /** Kind of device (GPU, CPU, or other). */
   cl_device_type type;
   /** Whether host memory is unified. */
@@ -349,6 +354,8 @@ typedef struct c_dbcsr_acc_opencl_config_t {
 /** Global configuration setup in c_dbcsr_acc_init. */
 extern c_dbcsr_acc_opencl_config_t c_dbcsr_acc_opencl_config;
 
+/** If buffers are hinted for non-concurrent writes aka "OpenCL constant". */
+int c_dbcsr_acc_opencl_use_cmem(const c_dbcsr_acc_opencl_device_t* devinfo);
 /** Determines host-pointer registration for modification. */
 c_dbcsr_acc_opencl_info_memptr_t* c_dbcsr_acc_opencl_info_hostptr(const void* memory);
 /** Determines device-pointer registration for modification (internal); offset is measured in elsize. */

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

@@ -70,7 +70,8 @@ __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, GLOBAL const int* restrict param_stack,
+FN(global T* restrict cdata, CONSTANT const T* restrict adata, CONSTANT const T* restrict bdata,
+  CONSTANT const int* restrict param_stack,
 #if (1 < BS)
   int param_format, int stack_size, int bs) {
   const int gid = get_group_id(0), idx = get_local_id(0);
@@ -80,11 +81,11 @@ FN(global T* restrict cdata, GLOBAL const T* restrict adata, GLOBAL const T* res
 #endif
   const SINT pzero = (0 == param_format ? 1 : 0), pnext = (0 == param_format ? 3 : 6);
   /* param_stack/indexes can be one-based (Fortran) depending on param_format */
-  GLOBAL const int* restrict param_base = param_stack + gid * (pnext * bs) + (0 == param_format ? 0 : 3);
+  CONSTANT const int* restrict param_base = param_stack + gid * (pnext * bs) + (0 == param_format ? 0 : 3);
 #if defined(SLM_P) && (1 < BS)
   local int params[3 * BS]; /* bs <= BS */
 #else
-  GLOBAL const int* restrict params = param_base;
+  CONSTANT const int* restrict params = param_base;
 #endif
 #if defined(SLM_A)
 #  if (1 != BK || BM < SM || 1 != BN)

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

@@ -8,7 +8,7 @@
 /*------------------------------------------------------------------------------------------------*/
 
 __attribute__((reqd_work_group_size(WG, 1, 1))) kernel void FN(
-  int trs_offset, GLOBAL const int* restrict trs_stack, global T* restrict matrix) {
+  int trs_offset, CONSTANT const int* restrict trs_stack, global T* restrict matrix) {
   /* offset in the transpose-stack that this block ID should handle */
   const int offset = trs_stack[trs_offset + get_group_id(0)];
   /* matrix according to the index (transpose-stack) */

src/acc/opencl/smm/opencl_libsmm.c

@@ -37,9 +37,6 @@
 #  if !defined(OPENCL_LIBSMM_NLOCKS_SMM)
 #    define OPENCL_LIBSMM_NLOCKS_SMM 16
 #  endif
-#  if !defined(OPENCL_LIBSMM_CMEM) && 1
-#    define OPENCL_LIBSMM_CMEM
-#  endif
 #  if !defined(OPENCL_LIBSMM_TODO) && 0
 #    define OPENCL_LIBSMM_TODO
 #  endif
@@ -90,22 +87,6 @@ opencl_libsmm_acc_dbm_launch_fn_t opencl_libsmm_acc_dbm_launch_fn;
 int opencl_libsmm_initialized;
 
 
-int opencl_libsmm_use_cmem(cl_device_id device) {
-#  if defined(OPENCL_LIBSMM_CMEM)
-  int result = EXIT_SUCCESS;
-  cl_ulong size_maxalloc = 1, size_maxcmem = 0;
-  ACC_OPENCL_CHECK(result, clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(cl_ulong), &size_maxalloc, NULL),
-    "retrieve maximum size of memory allocation");
-  ACC_OPENCL_CHECK(result, clGetDeviceInfo(device, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, sizeof(cl_ulong), &size_maxcmem, NULL),
-    "retrieve maximum size of constant buffer");
-  return (EXIT_SUCCESS == result ? (size_maxalloc <= size_maxcmem ? EXIT_SUCCESS : EXIT_FAILURE) : result);
-#  else
-  LIBXSMM_UNUSED(device);
-  return EXIT_FAILURE;
-#  endif
-}
-
-
 int opencl_libsmm_write_trans_params(FILE* stream, int only_key, const opencl_libsmm_transkey_t* key,
   const opencl_libsmm_trans_t* config, const char* delim, const char* begin, const char* close) {
   int result = 0;
@@ -699,8 +680,8 @@ int libsmm_acc_transpose(const int* dev_trs_stack, int offset, int stack_size, v
         const cl_device_id device_id = c_dbcsr_acc_opencl_config.devices[c_dbcsr_acc_opencl_config.device_id];
         const c_dbcsr_acc_opencl_device_t* const devinfo = &c_dbcsr_acc_opencl_config.device;
         const char *const env_cl = OPENCL_LIBSMM_TRANSENV("BUILDOPTS"), *const env_bm = OPENCL_LIBSMM_TRANSENV("BM");
-        const char* const cmem = (EXIT_SUCCESS != opencl_libsmm_use_cmem(device_id) ? "global" : "constant");
-        const char* const build_format = "-DGLOBAL=%s -DINPLACE=%i -DFN=%s -DSM=%i -DSN=%i -DWG=%i -DT=%s";
+        const char* const cmem = (EXIT_SUCCESS != c_dbcsr_acc_opencl_use_cmem(devinfo) ? "global" : "constant");
+        const char* const build_format = "-DCONSTANT=%s -DINPLACE=%i -DFN=%s -DSM=%i -DSN=%i -DWG=%i -DT=%s";
         const char *const env_inplace = OPENCL_LIBSMM_TRANSENV("INPLACE"), *tname = "";
 #  if defined(OPENCL_LIBSMM_TRANS_INPLACE)
         const int inplace = ((m == n) && (NULL == env_inplace ? 1 : ('0' != *env_inplace)));
@@ -1079,15 +1060,15 @@ int opencl_libsmm_acc_process(const int* host_param_stack, const int* dev_param_
               new_config.wgsize[kernel_idx] = (2 > new_config.wg ? (nbm * nbn) : ((int)LIBXSMM_UP2POT(nbm * nbn)));
             }
             if (new_config.wgsize[kernel_idx] <= devinfo->wgsize[0]) { /* SMM can be handled by device */
-              const char* const cmem = (EXIT_SUCCESS != opencl_libsmm_use_cmem(device_id) ? "global" : "constant");
+              const char* const cmem = (EXIT_SUCCESS != c_dbcsr_acc_opencl_use_cmem(devinfo) ? "global" : "constant");
               const char* const env_nrepeat = getenv("NREPEAT_SMM");
               const int typesize = OPENCL_LIBSMM_TYPESIZE(datatype);
               const int slm_a = (1 != new_config.aa ? 0 : (LIBXSMM_ISPOT(k_max * typesize) + 1));
               const int slm_b = (1 != new_config.ab ? 0 : (LIBXSMM_ISPOT(k_max * typesize) + 1));
               const int slm_c = (1 != new_config.ac ? 0 : (LIBXSMM_ISPOT(m_max * typesize) + 1));
               /* compose build parameters and flags */
               nchar = LIBXSMM_SNPRINTF(build_params, sizeof(build_params),
-                "-DT=%s -DGPU=%u -DGLOBAL=%s -DWG=%i -DSG=%i -DFN=%s -DREPEAT=%i -DLU=%i "
+                "-DT=%s -DGPU=%u -DCONSTANT=%s -DWG=%i -DSG=%i -DFN=%s -DREPEAT=%i -DLU=%i "
                 "-DSM=%i -DSN=%i -DSK=%i -DBS=%i -DVL=%i %s -DBM=%i -DBN=%i -DBK=%i "
                 "%s %s %s %s %s %s %s %s ", /* space! */
                 tname, CL_DEVICE_TYPE_GPU == devinfo->type, cmem, (int)new_config.wgsize[kernel_idx], (int)sgs, fname,

src/acc/opencl/smm/opencl_libsmm.h

@@ -74,9 +74,6 @@ typedef struct opencl_libsmm_perfest_t {
 /** Returns environment variable's value for given domain and key. */
 const char* opencl_libsmm_getenv(const char domain[], const char key[]);
 
-/** If buffers are hinted for non-concurrent writes aka "OpenCL constant". */
-int opencl_libsmm_use_cmem(cl_device_id device);
-
 /**
  * TRANS-kernel: write key and tunables into a (file-)stream.
  * If config=NULL, key/parameter names are written. The arguments