diff --git a/src/acc/opencl/Makefile b/src/acc/opencl/Makefile
index 244a7b2692f..211f9957113 100644
--- a/src/acc/opencl/Makefile
+++ b/src/acc/opencl/Makefile
@@ -22,10 +22,12 @@ endif
 UNAME := $(shell uname)
 HEADERONLY ?= 0
 STATIC ?= 1
-INTEL ?= 0
-GNU ?= 0
 DEV ?= 0
 
+# Intel Compiler
+ICX := $(shell which icx 2>/dev/null)
+INTEL ?= $(if $(ICX),$(if $(filter-out 0,$(GNU)),0,2),0)
+
 # select from set of predefined triplet specifications
 SPECID ?= 0
 # limit shape in tests (zero or negative for unlimited)
@@ -200,11 +202,15 @@ else
     endif
   else ifeq (,$(OPENCL_INC))
     ifneq (,$(wildcard $(OPENCL_ROOT)/include/CL/cl.h))
+      LDFLAGS += -L$(OPENCL_ROOT)/$(if $(wildcard $(OPENCL_ROOT)/lib64),lib64,lib)
       OPENCL_INC := $(OPENCL_ROOT)/include
-      LDFLAGS += -L$(OPENCL_ROOT)/lib64
-    else ifneq (,$(wildcard $(OPENCL_ROOT)/include/sycl/CL/cl.h))
-      OPENCL_INC := $(OPENCL_ROOT)/include/sycl
-      LDFLAGS += -L$(OPENCL_ROOT)/compiler/lib/intel64 -lintlc
+    else ifneq (,$(ICX))
+      OPENCL_ROOT := $(abspath $(dir $(ICX))/..)
+      ifneq (,$(wildcard $(OPENCL_ROOT)/include/sycl/CL/cl.h))
+        LDFLAGS += -L$(OPENCL_ROOT)/$(if $(wildcard $(OPENCL_ROOT)/lib64),lib64,lib)
+        LDFLAGS += -L$(OPENCL_ROOT)/compiler/lib/intel64 -lintlc
+        OPENCL_INC := $(OPENCL_ROOT)/include/sycl
+      endif
     endif
   endif
   # OPENCL_INC: directory containing CL/cl.h.
diff --git a/src/acc/opencl/acc_opencl.c b/src/acc/opencl/acc_opencl.c
index d7ee7651184..27e98070b51 100644
--- a/src/acc/opencl/acc_opencl.c
+++ b/src/acc/opencl/acc_opencl.c
@@ -58,6 +58,10 @@
 #  if !defined(ACC_OPENCL_SEDBIN) && 1
 #    define ACC_OPENCL_SEDBIN "/usr/bin/sed"
 #  endif
+/* disabled: let MPI runtime come up before */
+#  if !defined(ACC_OPENCL_PREINIT) && 0
+#    define ACC_OPENCL_PREINIT
+#  endif
 /* attempt to enable command aggregation */
 #  if !defined(ACC_OPENCL_CMDAGR) && 1
 #    define ACC_OPENCL_CMDAGR
@@ -151,87 +155,23 @@ int c_dbcsr_acc_opencl_order_devices(const void* dev_a, const void* dev_b) {
 }
 
 
-/* attempt to  automatically initialize backend */
-LIBXSMM_ATTRIBUTE_CTOR void c_dbcsr_acc_opencl_init(void) { ACC_OPENCL_EXPECT(EXIT_SUCCESS == c_dbcsr_acc_init()); }
-
-
-/* attempt to automatically finalize backend */
-LIBXSMM_ATTRIBUTE_DTOR void c_dbcsr_acc_opencl_finalize(void) {
-  assert(c_dbcsr_acc_opencl_config.ndevices < ACC_OPENCL_MAXNDEVS);
-  if (0 != c_dbcsr_acc_opencl_config.ndevices) {
-    int i;
-    for (i = 0; i < ACC_OPENCL_MAXNDEVS; ++i) {
-      const cl_device_id device_id = c_dbcsr_acc_opencl_config.devices[i];
-      if (NULL != device_id) {
-#  if defined(CL_VERSION_1_2) && 0 /* avoid potential segfault */
-        ACC_OPENCL_EXPECT(EXIT_SUCCESS == clReleaseDevice(device_id));
-#  endif
-        /* c_dbcsr_acc_opencl_create_context scans for non-NULL devices */
-        c_dbcsr_acc_opencl_config.devices[i] = NULL;
-      }
-    }
-    if (NULL != c_dbcsr_acc_opencl_config.device.stream.queue) { /* release private stream */
-      clReleaseCommandQueue(c_dbcsr_acc_opencl_config.device.stream.queue); /* ignore return code */
-    }
-    if (NULL != c_dbcsr_acc_opencl_config.device.context) {
-      const cl_context context = c_dbcsr_acc_opencl_config.device.context;
-      c_dbcsr_acc_opencl_config.device.context = NULL;
-      clReleaseContext(context); /* ignore return code */
-    }
-    for (i = 0; i < ACC_OPENCL_NLOCKS; ++i) { /* destroy locks */
-      ACC_OPENCL_DESTROY((ACC_OPENCL_LOCKTYPE*)(c_dbcsr_acc_opencl_locks + ACC_OPENCL_CACHELINE * i));
-    }
-    /* release/reset buffers */
-#  if defined(ACC_OPENCL_MEM_DEVPTR)
-    free(c_dbcsr_acc_opencl_config.memptrs);
-    free(c_dbcsr_acc_opencl_config.memptr_data);
-#  endif
-    free(c_dbcsr_acc_opencl_config.streams);
-    free(c_dbcsr_acc_opencl_config.stream_data);
-    free(c_dbcsr_acc_opencl_config.events);
-    free(c_dbcsr_acc_opencl_config.event_data);
-    /* clear entire configuration structure */
-    memset(&c_dbcsr_acc_opencl_config, 0, sizeof(c_dbcsr_acc_opencl_config));
-#  if defined(ACC_OPENCL_CACHE_DID)
-    c_dbcsr_acc_opencl_active_id = 0; /* reset cached active device-ID */
-#  endif
-    libxsmm_finalize();
-  }
-}
-
-
-int c_dbcsr_acc_init(void) {
-#  if defined(_OPENMP)
-  /* initialization/finalization is not meant to be thread-safe */
-  int result = ((0 == omp_in_parallel() || /*main*/ 0 == omp_get_thread_num()) ? EXIT_SUCCESS : EXIT_FAILURE);
-#  else
-  int result = EXIT_SUCCESS;
-#  endif
-#  if defined(__DBCSR_ACC) && defined(ACC_OPENCL_PROFILE)
-  int routine_handle;
-  static const char* const routine_name_ptr = LIBXSMM_FUNCNAME;
-  static const int routine_name_len = (int)sizeof(LIBXSMM_FUNCNAME) - 1;
-  c_dbcsr_timeset((const char**)&routine_name_ptr, &routine_name_len, &routine_handle);
-#  endif
-  if (EXIT_SUCCESS == result && 0 == c_dbcsr_acc_opencl_config.ndevices) { /* avoid to initialize multiple times */
-    cl_platform_id platforms[ACC_OPENCL_MAXNDEVS] = {NULL};
-    cl_device_id devices[ACC_OPENCL_MAXNDEVS];
-    char buffer[ACC_OPENCL_BUFFERSIZE];
-    const char *const env_devsplit = getenv("ACC_OPENCL_DEVSPLIT"), *const env_priority = getenv("ACC_OPENCL_PRIORITY");
-    const char *const env_devmatch = getenv("ACC_OPENCL_DEVMATCH"), *const env_devtype = getenv("ACC_OPENCL_DEVTYPE");
-    const char *const env_verbose = getenv("ACC_OPENCL_VERBOSE"), *const env_debug = getenv("ACC_OPENCL_DEBUG");
-    const char *const env_device = getenv("ACC_OPENCL_DEVICE"), *const env_dump_acc = getenv("ACC_OPENCL_DUMP");
+void c_dbcsr_acc_opencl_configure(void) {
+  if (NULL == c_dbcsr_acc_opencl_config.lock_main) { /* avoid to initialize multiple times */
+    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"), *const env_devsplit = getenv("ACC_OPENCL_DEVSPLIT");
+    const char *const env_verbose = getenv("ACC_OPENCL_VERBOSE"), *const env_dump_acc = getenv("ACC_OPENCL_DUMP");
+    const char *const env_priority = getenv("ACC_OPENCL_PRIORITY"), *const env_debug = getenv("ACC_OPENCL_DEBUG");
     const char *const env_timer = getenv("ACC_OPENCL_TIMER"), *const env_nlocks = getenv("ACC_OPENCL_NLOCKS");
     const char* const env_dump = (NULL != env_dump_acc ? env_dump_acc : getenv("IGC_ShaderDumpEnable"));
+    const char *const env_neo = getenv("NEOReadDebugKeys"), *const env_wa = getenv("ACC_OPENCL_WA");
+    static char neo_enable_debug_keys[] = "NEOReadDebugKeys=1";
 #  if defined(ACC_OPENCL_NCCS)
     const char* const env_nccs = getenv("ACC_OPENCL_NCCS");
     const int nccs = (NULL == env_nccs ? ACC_OPENCL_NCCS : atoi(env_nccs));
 #  endif
-    const char *const env_neo = getenv("NEOReadDebugKeys"), *const env_wa = getenv("ACC_OPENCL_WA");
-    const int neo = (NULL == env_neo ? 1 : atoi(env_neo));
 #  if defined(ACC_OPENCL_XHINTS)
     const char* const env_xhints = (ACC_OPENCL_XHINTS);
-    const int xhints_default = 1 + 2 + 4 + 8;
+    const int xhints_default = 1 + 2 + 4 + 8 + 16;
 #  else
     const char* const env_xhints = NULL;
     const int xhints_default = 0;
@@ -243,14 +183,9 @@ int c_dbcsr_acc_init(void) {
     const char* const env_async = NULL;
     const int async_default = 0;
 #  endif
-    /*const char* const env_nranks = getenv("MPI_LOCALNRANKS");
-    const cl_uint nranks = LIBXSMM_MAX(NULL != env_nranks ? atoi(env_nranks) : 1, 1);*/
-    const cl_int devsplit = (NULL == env_devsplit ? /*(1 < nranks ? -1 : 0)*/ 0 : atoi(env_devsplit));
-    char* const env_devids = getenv("ACC_OPENCL_DEVIDS");
-    int device_id = (NULL == env_device ? 0 : atoi(env_device));
     const int nlocks = (NULL == env_nlocks ? 1 /*default*/ : atoi(env_nlocks));
-    cl_uint nplatforms = 0, ndevices = 0, i;
-    cl_device_type type = CL_DEVICE_TYPE_ALL;
+    const int neo = (NULL == env_neo ? 1 : atoi(env_neo));
+    int i;
 #  if defined(_OPENMP)
     const int max_threads = omp_get_max_threads(), num_threads = omp_get_num_threads();
     c_dbcsr_acc_opencl_config.nthreads = (num_threads < max_threads ? max_threads : num_threads);
@@ -260,11 +195,9 @@ int c_dbcsr_acc_init(void) {
     c_dbcsr_acc_opencl_config.nthreads = 1;
     c_dbcsr_acc_opencl_config.nstreams = ACC_OPENCL_MAXNITEMS;
 #  endif
-#  if defined(ACC_OPENCL_CACHE_DID)
-    assert(0 == c_dbcsr_acc_opencl_active_id);
-#  endif
+    c_dbcsr_acc_opencl_config.nranks = LIBXSMM_MAX(NULL != env_nranks ? atoi(env_nranks) : 1, 1);
+    c_dbcsr_acc_opencl_config.nrank = (NULL != env_rank ? atoi(env_rank) : 0) % c_dbcsr_acc_opencl_config.nranks;
     assert(sizeof(ACC_OPENCL_LOCKTYPE) <= ACC_OPENCL_CACHELINE);
-    assert(EXIT_SUCCESS == result);
     for (i = 0; i < ACC_OPENCL_NLOCKS; ++i) {
       ACC_OPENCL_INIT((ACC_OPENCL_LOCKTYPE*)(c_dbcsr_acc_opencl_locks + ACC_OPENCL_CACHELINE * i));
     }
@@ -279,16 +212,15 @@ int c_dbcsr_acc_init(void) {
       (3 < LIBXSMM_MIN(nlocks, ACC_OPENCL_NLOCKS) ? ((ACC_OPENCL_LOCKTYPE*)(c_dbcsr_acc_opencl_locks + ACC_OPENCL_CACHELINE * 3))
                                                   : c_dbcsr_acc_opencl_config.lock_main);
     c_dbcsr_acc_opencl_config.verbosity = (NULL == env_verbose ? 0 : atoi(env_verbose));
+    c_dbcsr_acc_opencl_config.devsplit = (NULL == env_devsplit ? (/*1 < c_dbcsr_acc_opencl_config.nranks ? -1 :*/ 0)
+                                                               : atoi(env_devsplit));
     c_dbcsr_acc_opencl_config.priority = (NULL == env_priority ? /*default*/ 3 : atoi(env_priority));
     c_dbcsr_acc_opencl_config.xhints = (NULL == env_xhints ? xhints_default : atoi(env_xhints));
     c_dbcsr_acc_opencl_config.async = (NULL == env_async ? async_default : atoi(env_async));
     c_dbcsr_acc_opencl_config.dump = (NULL == env_dump ? /*default*/ 0 : atoi(env_dump));
     c_dbcsr_acc_opencl_config.debug = (NULL == env_debug ? c_dbcsr_acc_opencl_config.dump : atoi(env_debug));
-    c_dbcsr_acc_opencl_config.wa = neo * (NULL == env_wa ? ((1 != devsplit ? 0 : 4) + (8 + 16 + 32)) : atoi(env_wa));
-    assert(EXIT_SUCCESS == result);
-    if (EXIT_SUCCESS != c_dbcsr_acc_opencl_device_uid(NULL /*device*/, env_devmatch, &c_dbcsr_acc_opencl_config.devmatch)) {
-      c_dbcsr_acc_opencl_config.devmatch = 1;
-    }
+    c_dbcsr_acc_opencl_config.wa = neo * (NULL == env_wa ? ((1 != c_dbcsr_acc_opencl_config.devsplit ? 0 : 2) + (4 + 8 + 16))
+                                                         : atoi(env_wa));
     libxsmm_init();
     if (NULL != env_timer && (c_dbcsr_acc_opencl_timer_host == atoi(env_timer) ||
                                (env_timer == LIBXSMM_STRISTR(env_timer, "host") && 4 == strlen(env_timer)) ||
@@ -296,7 +228,6 @@ int c_dbcsr_acc_init(void) {
     {
       c_dbcsr_acc_opencl_config.timer = c_dbcsr_acc_opencl_timer_host;
     }
-    assert(EXIT_SUCCESS == result);
 #  if defined(ACC_OPENCL_CACHE_DIR)
     { /* environment is populated before touching the compute runtime */
       const char *const env_cache = getenv("ACC_OPENCL_CACHE"), *env_cachedir = getenv("NEO_CACHE_DIR");
@@ -339,10 +270,9 @@ int c_dbcsr_acc_init(void) {
 #    endif
       }
     }
-    assert(EXIT_SUCCESS == result);
 #  endif
 #  if defined(ACC_OPENCL_NCCS)
-    if ((1 & c_dbcsr_acc_opencl_config.wa) && 0 != nccs && NULL == getenv("ZEX_NUMBER_OF_CCS")) {
+    if (0 != nccs && NULL == getenv("ZEX_NUMBER_OF_CCS")) {
       static char zex_nccs[ACC_OPENCL_MAXNDEVS * 8 + 32] = "ZEX_NUMBER_OF_CCS=";
       const int mode = ((1 == nccs || 2 == nccs) ? nccs : 4);
       int j = strlen(zex_nccs);
@@ -361,10 +291,36 @@ int c_dbcsr_acc_init(void) {
         fprintf(stderr, "INFO ACC/OpenCL: support multiple separate compute command streamers (%i-CCS mode)\n", mode);
       }
     }
-    assert(EXIT_SUCCESS == result);
 #  endif
+    if (0 != neo && (NULL != env_neo || 0 == LIBXSMM_PUTENV(neo_enable_debug_keys))) {
+      if ((2 + 4 + 8) & c_dbcsr_acc_opencl_config.wa) {
+        static char a[] = "ZE_FLAT_DEVICE_HIERARCHY=COMPOSITE", b[] = "EnableRecoverablePageFaults=0";
+        static char c[] = "DirectSubmissionOverrideBlitterSupport=0", *const apply[] = {a, b, c};
+        if ((2 & c_dbcsr_acc_opencl_config.wa) && NULL == getenv("ZE_FLAT_DEVICE_HIERARCHY")) {
+          ACC_OPENCL_EXPECT(0 == LIBXSMM_PUTENV(apply[0]));
+        }
+        if ((4 & c_dbcsr_acc_opencl_config.wa) && NULL == getenv("EnableRecoverablePageFaults")) {
+          ACC_OPENCL_EXPECT(0 == LIBXSMM_PUTENV(apply[1]));
+        }
+        if ((8 & c_dbcsr_acc_opencl_config.wa) && NULL == getenv("DirectSubmissionOverrideBlitterSupport")) {
+          ACC_OPENCL_EXPECT(0 == LIBXSMM_PUTENV(apply[2]));
+        }
+      }
+      if (0 != c_dbcsr_acc_opencl_config.debug) {
+        static char a[] = "DisableScratchPages=1", *const apply[] = {a};
+        if (NULL == getenv("DisableScratchPages")) {
+          ACC_OPENCL_EXPECT(0 == LIBXSMM_PUTENV(apply[0]));
+        }
+      }
+    }
+    if (2 & c_dbcsr_acc_opencl_config.xhints) {
+      static char a[] = "I_MPI_OFFLOAD_RDMA=1", *const apply[] = {a};
+      if (NULL == getenv("I_MPI_OFFLOAD_RDMA")) {
+        ACC_OPENCL_EXPECT(0 == LIBXSMM_PUTENV(apply[0]));
+      }
+    }
 #  if defined(ACC_OPENCL_DL)
-    if (2 & c_dbcsr_acc_opencl_config.wa) { /* initialize L0 */
+    if (1 & c_dbcsr_acc_opencl_config.wa) { /* initialize L0 */
       union {
         const void* dlsym;
         int (*ptr)(int flags);
@@ -381,23 +337,43 @@ int c_dbcsr_acc_init(void) {
         }
       }
     }
-    assert(EXIT_SUCCESS == result);
-#  endif
-    if (~(1 + 2 + 32) & c_dbcsr_acc_opencl_config.wa) { /* environment is populated before touching the compute runtime */
-      static char a[] = "NEOReadDebugKeys=1", b[] = "ZE_FLAT_DEVICE_HIERARCHY=COMPOSITE", c[] = "EnableRecoverablePageFaults=0";
-      static char d[] = "DirectSubmissionOverrideBlitterSupport=0", *key_value[] = {a, b, c, d};
-      if (NULL == env_neo) ACC_OPENCL_EXPECT(0 == LIBXSMM_PUTENV(key_value[0]));
-      if ((4 & c_dbcsr_acc_opencl_config.wa) && NULL == getenv("ZE_FLAT_DEVICE_HIERARCHY")) {
-        ACC_OPENCL_EXPECT(0 == LIBXSMM_PUTENV(key_value[1]));
-      }
-      if ((8 & c_dbcsr_acc_opencl_config.wa) && NULL == getenv("EnableRecoverablePageFaults")) {
-        ACC_OPENCL_EXPECT(0 == LIBXSMM_PUTENV(key_value[2]));
-      }
-      if ((16 & c_dbcsr_acc_opencl_config.wa) && NULL == getenv("DirectSubmissionOverrideBlitterSupport")) {
-        ACC_OPENCL_EXPECT(0 == LIBXSMM_PUTENV(key_value[3]));
-      }
+#  endif
+  }
+}
+
+
+int c_dbcsr_acc_init(void) {
+#  if defined(_OPENMP)
+  /* initialization/finalization is not meant to be thread-safe */
+  int result = ((0 == omp_in_parallel() || /*main*/ 0 == omp_get_thread_num()) ? EXIT_SUCCESS : EXIT_FAILURE);
+#  else
+  int result = EXIT_SUCCESS;
+#  endif
+#  if defined(__DBCSR_ACC) && defined(ACC_OPENCL_PROFILE)
+  int routine_handle;
+  static const char* const routine_name_ptr = LIBXSMM_FUNCNAME;
+  static const int routine_name_len = (int)sizeof(LIBXSMM_FUNCNAME) - 1;
+  c_dbcsr_timeset((const char**)&routine_name_ptr, &routine_name_len, &routine_handle);
+#  endif
+  /* eventually touch OpenCL/compute runtime after configure */
+  c_dbcsr_acc_opencl_configure();
+  if (0 == c_dbcsr_acc_opencl_config.ndevices && EXIT_SUCCESS == result) { /* avoid to initialize multiple times */
+    char buffer[ACC_OPENCL_BUFFERSIZE];
+    cl_platform_id platforms[ACC_OPENCL_MAXNDEVS] = {NULL};
+    cl_device_id devices[ACC_OPENCL_MAXNDEVS];
+    cl_device_type type = CL_DEVICE_TYPE_ALL;
+    cl_uint nplatforms = 0, ndevices = 0, i;
+    const char* const env_devmatch = getenv("ACC_OPENCL_DEVMATCH");
+    const char* const env_devtype = getenv("ACC_OPENCL_DEVTYPE");
+    const char* const env_device = getenv("ACC_OPENCL_DEVICE");
+    char* const env_devids = getenv("ACC_OPENCL_DEVIDS");
+    int device_id = (NULL == env_device ? 0 : atoi(env_device));
+#  if defined(ACC_OPENCL_CACHE_DID)
+    assert(0 == c_dbcsr_acc_opencl_active_id);
+#  endif
+    if (EXIT_SUCCESS != c_dbcsr_acc_opencl_device_uid(NULL /*device*/, env_devmatch, &c_dbcsr_acc_opencl_config.devmatch)) {
+      c_dbcsr_acc_opencl_config.devmatch = 1;
     }
-    assert(EXIT_SUCCESS == result);
     if (EXIT_SUCCESS == clGetPlatformIDs(0, NULL, &nplatforms) && 0 < nplatforms) {
       ACC_OPENCL_CHECK(clGetPlatformIDs(nplatforms <= ACC_OPENCL_MAXNDEVS ? nplatforms : ACC_OPENCL_MAXNDEVS, platforms, 0),
         "retrieve platform ids", result);
@@ -428,15 +404,18 @@ int c_dbcsr_acc_init(void) {
               cl_device_partition_property properties[] = {
                 CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN, CL_DEVICE_AFFINITY_DOMAIN_NUMA, /*terminator*/ 0};
               cl_uint nunits = 0, n = 0;
-              if ((1 < devsplit || 0 > devsplit) && /* Intel CPU (e.g., out of two sockets) yields thread-count of both sockets */
+              if ((1 < c_dbcsr_acc_opencl_config.devsplit || 0 > c_dbcsr_acc_opencl_config.devsplit) &&
+                  /* Intel CPU (e.g., out of two sockets) yields thread-count of both sockets */
                   EXIT_SUCCESS == clGetDeviceInfo(devices[j], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint), &nunits, NULL) &&
                   1 < nunits)
               {
-                n = LIBXSMM_MIN(1 < devsplit ? (cl_uint)devsplit : nunits, ACC_OPENCL_MAXNDEVS);
+                n = LIBXSMM_MIN(1 < c_dbcsr_acc_opencl_config.devsplit ? (cl_uint)c_dbcsr_acc_opencl_config.devsplit : nunits,
+                  ACC_OPENCL_MAXNDEVS);
                 properties[0] = CL_DEVICE_PARTITION_EQUALLY;
                 properties[1] = (nunits + n - 1) / n;
               }
-              if (0 == devsplit || 1 == devsplit || (c_dbcsr_acc_opencl_config.ndevices + 1) == ACC_OPENCL_MAXNDEVS ||
+              if (0 == c_dbcsr_acc_opencl_config.devsplit || 1 == c_dbcsr_acc_opencl_config.devsplit ||
+                  (c_dbcsr_acc_opencl_config.ndevices + 1) == ACC_OPENCL_MAXNDEVS ||
                   EXIT_SUCCESS != clCreateSubDevices(devices[j], properties, 0, NULL, &n))
 #  endif
               {
@@ -651,17 +630,17 @@ int c_dbcsr_acc_init(void) {
             result = c_dbcsr_acc_opencl_set_active_device(NULL /*lock*/, ACC_OPENCL_ACTIVATE);
           }
           else {
-            const char* const env_rank = (NULL != getenv("PMI_RANK") ? getenv("PMI_RANK") : getenv("OMPI_COMM_WORLD_LOCAL_RANK"));
-            const int rank = (NULL != env_rank ? atoi(env_rank) : 0);
-            if (0 < rank && 1 < c_dbcsr_acc_opencl_config.ndevices) {
-              device_id = rank % c_dbcsr_acc_opencl_config.ndevices;
+            if (0 < c_dbcsr_acc_opencl_config.nrank && 1 < c_dbcsr_acc_opencl_config.ndevices) {
+              device_id = c_dbcsr_acc_opencl_config.nrank % c_dbcsr_acc_opencl_config.ndevices;
             }
             result = c_dbcsr_acc_opencl_set_active_device(NULL /*lock*/, device_id);
           }
 #  else
           c_dbcsr_acc_opencl_config.device.uid = (cl_uint)device_id; /* hack */
 #  endif
-          if (2 <= c_dbcsr_acc_opencl_config.verbosity || 0 > c_dbcsr_acc_opencl_config.verbosity) {
+          if ((2 <= c_dbcsr_acc_opencl_config.verbosity || 0 > c_dbcsr_acc_opencl_config.verbosity) &&
+              (0 == c_dbcsr_acc_opencl_config.nrank))
+          {
             char platform_name[ACC_OPENCL_BUFFERSIZE];
             for (i = 0; i < (cl_uint)c_dbcsr_acc_opencl_config.ndevices; ++i) {
               if (EXIT_SUCCESS == c_dbcsr_acc_opencl_device_name(c_dbcsr_acc_opencl_config.devices[i], buffer,
@@ -693,6 +672,58 @@ int c_dbcsr_acc_init(void) {
 }
 
 
+/* attempt to automatically initialize backend */
+LIBXSMM_ATTRIBUTE_CTOR void c_dbcsr_acc_opencl_init(void) {
+  c_dbcsr_acc_opencl_configure();
+#  if defined(ACC_OPENCL_PREINIT)
+  ACC_OPENCL_EXPECT(EXIT_SUCCESS == c_dbcsr_acc_init());
+#  endif
+}
+
+
+/* attempt to automatically finalize backend */
+LIBXSMM_ATTRIBUTE_DTOR void c_dbcsr_acc_opencl_finalize(void) {
+  assert(c_dbcsr_acc_opencl_config.ndevices < ACC_OPENCL_MAXNDEVS);
+  if (0 != c_dbcsr_acc_opencl_config.ndevices) {
+    int i;
+    for (i = 0; i < ACC_OPENCL_MAXNDEVS; ++i) {
+      const cl_device_id device_id = c_dbcsr_acc_opencl_config.devices[i];
+      if (NULL != device_id) {
+#  if defined(CL_VERSION_1_2) && 0 /* avoid potential segfault */
+        ACC_OPENCL_EXPECT(EXIT_SUCCESS == clReleaseDevice(device_id));
+#  endif
+      }
+    }
+    if (NULL != c_dbcsr_acc_opencl_config.device.stream.queue) { /* release private stream */
+      clReleaseCommandQueue(c_dbcsr_acc_opencl_config.device.stream.queue); /* ignore return code */
+    }
+    if (NULL != c_dbcsr_acc_opencl_config.device.context) {
+      const cl_context context = c_dbcsr_acc_opencl_config.device.context;
+      c_dbcsr_acc_opencl_config.device.context = NULL;
+      clReleaseContext(context); /* ignore return code */
+    }
+    for (i = 0; i < ACC_OPENCL_NLOCKS; ++i) { /* destroy locks */
+      ACC_OPENCL_DESTROY((ACC_OPENCL_LOCKTYPE*)(c_dbcsr_acc_opencl_locks + ACC_OPENCL_CACHELINE * i));
+    }
+    /* release/reset buffers */
+#  if defined(ACC_OPENCL_MEM_DEVPTR)
+    free(c_dbcsr_acc_opencl_config.memptrs);
+    free(c_dbcsr_acc_opencl_config.memptr_data);
+#  endif
+    free(c_dbcsr_acc_opencl_config.streams);
+    free(c_dbcsr_acc_opencl_config.stream_data);
+    free(c_dbcsr_acc_opencl_config.events);
+    free(c_dbcsr_acc_opencl_config.event_data);
+    /* clear entire configuration structure */
+    memset(&c_dbcsr_acc_opencl_config, 0, sizeof(c_dbcsr_acc_opencl_config));
+#  if defined(ACC_OPENCL_CACHE_DID)
+    c_dbcsr_acc_opencl_active_id = 0; /* reset cached active device-ID */
+#  endif
+    libxsmm_finalize();
+  }
+}
+
+
 int c_dbcsr_acc_finalize(void) {
 #  if defined(_OPENMP)
   /* initialization/finalization is not meant to be thread-safe */
diff --git a/src/acc/opencl/acc_opencl.h b/src/acc/opencl/acc_opencl.h
index ddd1834051f..41b27db6def 100644
--- a/src/acc/opencl/acc_opencl.h
+++ b/src/acc/opencl/acc_opencl.h
@@ -338,8 +338,12 @@ typedef struct c_dbcsr_acc_opencl_config_t {
   c_dbcsr_acc_opencl_timer_t timer; /* c_dbcsr_acc_opencl_device_t? */
   /** Kernel-parameters are matched against device's UID */
   cl_uint devmatch;
+  /** Split devices into sub-devices (if possible) */
+  cl_int devsplit;
   /** Verbosity level (output on stderr). */
   cl_int verbosity;
+  /** Guessed number of ranks per node (local), and rank-ID. */
+  cl_int nranks, nrank;
   /** Non-zero if library is initialized (negative: no device). */
   cl_int ndevices;
   /** Maximum number of threads (omp_get_max_threads). */
diff --git a/src/acc/opencl/acc_opencl_mem.c b/src/acc/opencl/acc_opencl_mem.c
index 7d9bd86a4d0..cc2d1a1b03c 100644
--- a/src/acc/opencl/acc_opencl_mem.c
+++ b/src/acc/opencl/acc_opencl_mem.c
@@ -189,7 +189,7 @@ int c_dbcsr_acc_host_mem_allocate(void** host_mem, size_t nbytes, void* stream)
   }
 #  endif
 #  if defined(ACC_OPENCL_XHINTS)
-  if (0 != (8 & c_dbcsr_acc_opencl_config.xhints) && (0 != c_dbcsr_acc_opencl_config.device.nv || NULL != (ACC_OPENCL_XHINTS))) {
+  if (0 != (16 & c_dbcsr_acc_opencl_config.xhints) && (0 != c_dbcsr_acc_opencl_config.device.nv || NULL != (ACC_OPENCL_XHINTS))) {
     host_ptr = malloc(nbytes);
     if (NULL != host_ptr) flags = CL_MEM_USE_HOST_PTR;
   }
@@ -202,7 +202,7 @@ int c_dbcsr_acc_host_mem_allocate(void** host_mem, size_t nbytes, void* stream)
                                                                      : c_dbcsr_acc_opencl_stream_default());
       mapped = clEnqueueMapBuffer(str->queue, memory, CL_TRUE /*always block*/,
 #  if defined(ACC_OPENCL_XHINTS) && (defined(CL_VERSION_1_2) || defined(CL_MAP_WRITE_INVALIDATE_REGION))
-        (4 & c_dbcsr_acc_opencl_config.xhints) ? CL_MAP_WRITE_INVALIDATE_REGION :
+        (8 & c_dbcsr_acc_opencl_config.xhints) ? CL_MAP_WRITE_INVALIDATE_REGION :
 #  endif
                                                (CL_MAP_READ | CL_MAP_WRITE),
         0 /*offset*/, nbytes, 0, NULL, NULL, &result);
@@ -250,7 +250,7 @@ int c_dbcsr_acc_host_mem_deallocate(void* host_mem, void* stream) {
       void* host_ptr = NULL;
       int result_release;
 #  if defined(ACC_OPENCL_XHINTS)
-      if (0 != (8 & c_dbcsr_acc_opencl_config.xhints) &&
+      if (0 != (16 & c_dbcsr_acc_opencl_config.xhints) &&
           (0 != c_dbcsr_acc_opencl_config.device.nv || NULL != (ACC_OPENCL_XHINTS)) &&
           EXIT_SUCCESS == clGetMemObjectInfo(info.memory, CL_MEM_HOST_PTR, sizeof(void*), &host_ptr, NULL) && NULL != host_ptr)
       {
@@ -343,79 +343,88 @@ int c_dbcsr_acc_dev_mem_allocate(void** dev_mem, size_t nbytes) {
   }
 #  endif
   assert(NULL != dev_mem && NULL != context);
+  if (0 != nbytes) {
 #  if defined(ACC_OPENCL_MEM_DEVPTR)
-  if (NULL != c_dbcsr_acc_opencl_config.device.clDeviceMemAllocINTEL) {
-    assert(0 == c_dbcsr_acc_opencl_config.device.unified);
-    *dev_mem = memptr = c_dbcsr_acc_opencl_config.device.clDeviceMemAllocINTEL(
-      context, c_dbcsr_acc_opencl_config.device.id, NULL /*properties*/, nbytes, 0 /*alignment*/, &result);
-    if (EXIT_SUCCESS != result) *dev_mem = NULL;
-  }
-  else
-#  endif
-  {
-    const int devuid = c_dbcsr_acc_opencl_config.device.uid;
-    const int try_flag = ((0 != c_dbcsr_acc_opencl_config.device.unified || 0 == c_dbcsr_acc_opencl_config.device.intel ||
-                            (0x4905 != devuid && 0x020a != devuid && (0x0bd0 > devuid || 0x0bdb < devuid)))
-                            ? 0
-                            : (1u << 22));
-    memory = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE | try_flag), nbytes, NULL /*host_ptr*/, &result);
-    if (0 != try_flag && EXIT_SUCCESS != result) { /* retry without try_flag */
-      memory = clCreateBuffer(context, CL_MEM_READ_WRITE, nbytes, NULL /*host_ptr*/, &result);
+    if (NULL != c_dbcsr_acc_opencl_config.device.clDeviceMemAllocINTEL) {
+      assert(0 == c_dbcsr_acc_opencl_config.device.unified);
+      *dev_mem = memptr = c_dbcsr_acc_opencl_config.device.clDeviceMemAllocINTEL(
+        context, c_dbcsr_acc_opencl_config.device.id, NULL /*properties*/, nbytes, 0 /*alignment*/, &result);
+      if (EXIT_SUCCESS != result) *dev_mem = NULL;
     }
-    if (EXIT_SUCCESS == result) {
-#  if defined(ACC_OPENCL_MEM_DEVPTR)
-      const c_dbcsr_acc_opencl_stream_t* str = NULL;
-      static cl_kernel kernel = NULL;
-      const size_t size = 1;
-      ACC_OPENCL_ACQUIRE(c_dbcsr_acc_opencl_config.lock_memory);
-      str = c_dbcsr_acc_opencl_stream(NULL /*lock*/, ACC_OPENCL_OMP_TID());
-      /* determine device-side value of device-memory object by running some kernel */
-      assert(NULL != memory && NULL != str && NULL != str->queue);
-      if (NULL == kernel) { /* generate kernel */
-        const char source[] = "kernel void memptr(global unsigned long* ptr) {\n"
-                              "  const union { global unsigned long* p; unsigned long u; } cast = { ptr };\n"
-                              "  const size_t i = get_global_id(0);\n"
-                              "  ptr[i] = cast.u + i;\n"
-                              "}\n";
-        assert(sizeof(size_t) == sizeof(cl_ulong));
-        result = c_dbcsr_acc_opencl_kernel(0 /*source_is_file*/, source, "memptr" /*kernel_name*/, NULL /*build_params*/,
-          NULL /*build_options*/, NULL /*try_build_options*/, NULL /*try_ok*/, NULL /*extnames*/, 0 /*num_exts*/, &kernel);
-      }
-      /* TODO: backup/restore memory */
-      if (EXIT_SUCCESS == result) result = clSetKernelArg(kernel, 0, sizeof(cl_mem), &memory);
-      if (EXIT_SUCCESS == result) {
-        result = clEnqueueNDRangeKernel(
-          str->queue, kernel, 1 /*work_dim*/, NULL /*offset*/, &size, NULL /*local_work_size*/, 0, NULL, NULL);
-      }
-      if (EXIT_SUCCESS == result) {
-        result = c_dbcsr_acc_opencl_memcpy_d2h(memory, &memptr, 0, sizeof(void*), str->queue, 1 /*blocking*/);
+    else
+#  endif
+    {
+      const int devuid = c_dbcsr_acc_opencl_config.device.uid;
+      const int try_flag = ((0 != c_dbcsr_acc_opencl_config.device.unified || 0 == c_dbcsr_acc_opencl_config.device.intel ||
+                              (0x4905 != devuid && 0x020a != devuid && (0x0bd0 > devuid || 0x0bdb < devuid)))
+                              ? 0
+                              : (1u << 22));
+      memory = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE | try_flag), nbytes, NULL /*host_ptr*/, &result);
+      if (0 != try_flag && EXIT_SUCCESS != result) { /* retry without try_flag */
+        memory = clCreateBuffer(context, CL_MEM_READ_WRITE, nbytes, NULL /*host_ptr*/, &result);
       }
-      assert(EXIT_SUCCESS != result || NULL != memptr);
       if (EXIT_SUCCESS == result) {
-        c_dbcsr_acc_opencl_info_memptr_t* const info = (c_dbcsr_acc_opencl_info_memptr_t*)c_dbcsr_acc_opencl_pmalloc(
-          NULL /*lock*/, (void**)c_dbcsr_acc_opencl_config.memptrs, &c_dbcsr_acc_opencl_config.nmemptrs);
-        assert(NULL != memptr);
-        if (NULL != info) {
-          info->memory = memory;
-          info->memptr = memptr;
-          *dev_mem = memptr;
+#  if defined(ACC_OPENCL_MEM_DEVPTR)
+        const c_dbcsr_acc_opencl_stream_t* str = NULL;
+        static cl_kernel kernel = NULL;
+        const size_t size = 1;
+        ACC_OPENCL_ACQUIRE(c_dbcsr_acc_opencl_config.lock_memory);
+        str = c_dbcsr_acc_opencl_stream(NULL /*lock*/, ACC_OPENCL_OMP_TID());
+        /* determine device-side value of device-memory object by running some kernel */
+        assert(NULL != memory && NULL != str && NULL != str->queue);
+        if (NULL == kernel) { /* generate kernel */
+          const char source[] = "kernel void memptr(global unsigned long* ptr) {\n"
+                                "  const union { global unsigned long* p; unsigned long u; } cast = { ptr };\n"
+                                "  const size_t i = get_global_id(0);\n"
+                                "  ptr[i] = cast.u + i;\n"
+                                "}\n";
+          assert(sizeof(size_t) == sizeof(cl_ulong));
+          result = c_dbcsr_acc_opencl_kernel(0 /*source_is_file*/, source, "memptr" /*kernel_name*/, NULL /*build_params*/,
+            NULL /*build_options*/, NULL /*try_build_options*/, NULL /*try_ok*/, NULL /*extnames*/, 0 /*num_exts*/, &kernel);
         }
-        else result = EXIT_FAILURE;
-      }
-      ACC_OPENCL_RELEASE(c_dbcsr_acc_opencl_config.lock_memory);
+        /* TODO: backup/restore memory */
+        if (EXIT_SUCCESS == result) result = clSetKernelArg(kernel, 0, sizeof(cl_mem), &memory);
+        if (EXIT_SUCCESS == result) {
+          result = clEnqueueNDRangeKernel(
+            str->queue, kernel, 1 /*work_dim*/, NULL /*offset*/, &size, NULL /*local_work_size*/, 0, NULL, NULL);
+        }
+        if (EXIT_SUCCESS == result) {
+          result = c_dbcsr_acc_opencl_memcpy_d2h(memory, &memptr, 0, sizeof(void*), str->queue, 1 /*blocking*/);
+        }
+        assert(EXIT_SUCCESS != result || NULL != memptr);
+        if (EXIT_SUCCESS == result) {
+          c_dbcsr_acc_opencl_info_memptr_t* const info = (c_dbcsr_acc_opencl_info_memptr_t*)c_dbcsr_acc_opencl_pmalloc(
+            NULL /*lock*/, (void**)c_dbcsr_acc_opencl_config.memptrs, &c_dbcsr_acc_opencl_config.nmemptrs);
+          assert(NULL != memptr);
+          if (NULL != info) {
+            info->memory = memory;
+            info->memptr = memptr;
+            *dev_mem = memptr;
+          }
+          else result = EXIT_FAILURE;
+        }
+        ACC_OPENCL_RELEASE(c_dbcsr_acc_opencl_config.lock_memory);
 #  else
-      *dev_mem = memptr = memory;
+        *dev_mem = memptr = memory;
 #  endif
+      }
+      if (EXIT_SUCCESS != result) {
+        if (NULL != memory) ACC_OPENCL_EXPECT(EXIT_SUCCESS == clReleaseMemObject(memory));
+        *dev_mem = NULL;
+      }
     }
-    if (EXIT_SUCCESS != result) {
-      if (NULL != memory) ACC_OPENCL_EXPECT(EXIT_SUCCESS == clReleaseMemObject(memory));
-      *dev_mem = NULL;
+    if (0 != c_dbcsr_acc_opencl_config.verbosity) {
+      if (EXIT_SUCCESS == result && 0 != c_dbcsr_acc_opencl_config.debug) {
+        fprintf(stderr, "INFO ACC/OpenCL: memory=%p pointer=%p size=%llu successfully allocated\n", (const void*)memory, memptr,
+          (unsigned long long)nbytes);
+      }
+      else if (EXIT_SUCCESS != result) {
+        fprintf(stderr, "ERROR ACC/OpenCL: memory=%p pointer=%p size=%llu failed to allocate\n", (const void*)memory, memptr,
+          (unsigned long long)nbytes);
+      }
     }
   }
-  if (0 != c_dbcsr_acc_opencl_config.debug && 0 != c_dbcsr_acc_opencl_config.verbosity && EXIT_SUCCESS == result) {
-    fprintf(stderr, "INFO ACC/OpenCL: memory=%p pointer=%p size=%llu allocated\n", (const void*)memory, memptr,
-      (unsigned long long)nbytes);
-  }
+  else *dev_mem = NULL;
   assert(EXIT_SUCCESS == result || NULL == *dev_mem);
 #  if defined(__DBCSR_ACC) && defined(ACC_OPENCL_PROFILE)
   c_dbcsr_timestop(&routine_handle);
diff --git a/src/acc/opencl/acc_opencl_stream.c b/src/acc/opencl/acc_opencl_stream.c
index 29ade32dba8..ce3b7e6aef7 100644
--- a/src/acc/opencl/acc_opencl_stream.c
+++ b/src/acc/opencl/acc_opencl_stream.c
@@ -118,7 +118,7 @@ int c_dbcsr_acc_stream_create(void** stream_p, const char* name, int priority) {
 #  endif
   {
 #  if defined(ACC_OPENCL_XHINTS)
-    if ((2 & c_dbcsr_acc_opencl_config.xhints) && 0 != c_dbcsr_acc_opencl_config.device.intel) { /* enable queue families */
+    if ((4 & c_dbcsr_acc_opencl_config.xhints) && 0 != c_dbcsr_acc_opencl_config.device.intel) { /* enable queue families */
       struct {
         cl_command_queue_properties properties;
         cl_bitfield capabilities;
diff --git a/src/acc/opencl/smm/opencl_libsmm.c b/src/acc/opencl/smm/opencl_libsmm.c
index 409659980cb..f5dc95a8683 100644
--- a/src/acc/opencl/smm/opencl_libsmm.c
+++ b/src/acc/opencl/smm/opencl_libsmm.c
@@ -451,6 +451,7 @@ int libsmm_acc_init(void) {
         }
 #  if defined(OPENCL_KERNELS_PARAMS_SMM) && defined(OPENCL_KERNELS_DEVICES)
         if (EXIT_SUCCESS == result && (0 == ntuned || 0 != key_direct_skip)) {
+          unsigned int default_uid = c_dbcsr_acc_opencl_config.device.uid;
           const char *line = OPENCL_KERNELS_PARAMS_SMM, *next;
 #    if LIBXSMM_VERSION4(1, 17, 0, 0) < LIBXSMM_VERSION_NUMBER
           int active_match = -1;
@@ -458,13 +459,16 @@ int libsmm_acc_init(void) {
                                 NULL /*platform*/, 0 /*platform_maxlen*/, /*cleanup*/ 1))
           { /* determine best-matching parameters based on name of device */
             int i = 0, best = 0;
+            if (1 >= c_dbcsr_acc_opencl_config.devmatch) {
+              c_dbcsr_acc_opencl_device_uid(c_dbcsr_acc_opencl_config.device.id, bufname, &default_uid);
+            }
             for (; i < ndevices_params; ++i) {
               const int score = libxsmm_strimatch(bufname, OPENCL_KERNELS_DEVICES[i], NULL);
               unsigned int uid;
               if (best < score ||
                   ((best == score) &&
                     EXIT_SUCCESS == c_dbcsr_acc_opencl_device_uid(NULL /*device*/, OPENCL_KERNELS_DEVICES[i], &uid) &&
-                    uid == c_dbcsr_acc_opencl_config.device.uid))
+                    uid == default_uid))
               {
                 active_match = i;
                 best = score;
@@ -489,22 +493,20 @@ int libsmm_acc_init(void) {
                   {
                     key.devuid = 0;
                   }
-                  config_init = (opencl_libsmm_smm_t*)libxsmm_xdispatch(&key, sizeof(key));
+                  config_init = (opencl_libsmm_smm_t*)libxsmm_xdispatch(&key, sizeof(key)); /* duplicate? */
                   if (NULL == config_init) {
-                    if (NULL == libxsmm_xregister(&key, sizeof(key), sizeof(config), &config)) {
+                    if (NULL != libxsmm_xregister(&key, sizeof(key), sizeof(config), &config)) ++ntuned;
+                    else { /* failed to register */
                       result = EXIT_FAILURE;
                       break;
                     }
-                    else ++ntuned;
                   }
                   else if (config_init->gflops < config.gflops) { /* update */
                     memcpy(config_init, &config, sizeof(config));
                   }
 #    if LIBXSMM_VERSION4(1, 17, 0, 0) < LIBXSMM_VERSION_NUMBER
-                  if (active_match == i && 0 != c_dbcsr_acc_opencl_config.device.uid &&
-                      c_dbcsr_acc_opencl_config.device.uid != key.devuid)
-                  {
-                    key.devuid = c_dbcsr_acc_opencl_config.device.uid;
+                  if (active_match == i && 0 != default_uid && default_uid != key.devuid) {
+                    key.devuid = default_uid;
                     config_init = (opencl_libsmm_smm_t*)libxsmm_xdispatch(&key, sizeof(key));
                     if (NULL == config_init && NULL != libxsmm_xregister(&key, sizeof(key), sizeof(config), &config)) {
                       static int info = 0;
@@ -532,7 +534,7 @@ int libsmm_acc_init(void) {
         }
 #  endif
 #  if defined(OPENCL_KERNELS_DEVICES)
-        if (EXIT_SUCCESS == result && 0 != ntuned &&
+        if (EXIT_SUCCESS == result && 0 != ntuned && 0 == c_dbcsr_acc_opencl_config.nrank &&
             (2 <= c_dbcsr_acc_opencl_config.verbosity || 0 > c_dbcsr_acc_opencl_config.verbosity))
         {
           fprintf(stderr, "INFO ACC/LIBSMM: PARAMS in %i set%s loaded targeting ", ntuned, 1 != ntuned ? "s" : "");
@@ -1084,7 +1086,7 @@ int libsmm_acc_process(const int* host_param_stack, const int* dev_param_stack,
               (NULL == env_nz || '\0' == *env_nz) ? (0 != defaults ? /*default*/ 0 : config->nz) : atoi(env_nz), 0, 1);
             new_config.al = LIBXSMM_CLMP(/* bug: AL=1 */
               (NULL == env_al || '\0' == *env_al)
-                ? (0 == (64 & c_dbcsr_acc_opencl_config.wa) ? (0 != defaults ? 0 : config->al) : 0)
+                ? (0 == (32 & c_dbcsr_acc_opencl_config.wa) ? (0 != defaults ? 0 : config->al) : 0)
                 : atoi(env_al),
               0, 1);
             new_config.tb = LIBXSMM_CLMP(
@@ -1095,7 +1097,7 @@ int libsmm_acc_process(const int* host_param_stack, const int* dev_param_stack,
               (NULL == env_ap || '\0' == *env_ap) ? (0 != defaults ? /*default*/ 0 : config->ap) : atoi(env_ap), 0, 1);
             new_config.aa = LIBXSMM_CLMP(/* bug: AA=2 XF=1 */
               (NULL == env_aa || '\0' == *env_aa) ? (0 != defaults ? default_aa : config->aa) : atoi(env_aa), 0,
-              (0 == (32 & c_dbcsr_acc_opencl_config.wa) || 0 == new_config.flags) ? 2 : 1);
+              (0 == (16 & c_dbcsr_acc_opencl_config.wa) || 0 == new_config.flags) ? 2 : 1);
             new_config.ab = LIBXSMM_CLMP(
               (NULL == env_ab || '\0' == *env_ab) ? (0 != defaults ? default_ab : config->ab) : atoi(env_ab), 0, 2);
             new_config.ac = LIBXSMM_CLMP(
diff --git a/src/acc/opencl/smm/params/tune_multiply_Mi250.csv b/src/acc/opencl/smm/params/tune_multiply_Mi250.csv
index d73036e2ca9..8067550ffe9 100644
--- a/src/acc/opencl/smm/params/tune_multiply_Mi250.csv
+++ b/src/acc/opencl/smm/params/tune_multiply_Mi250.csv
@@ -178,7 +178,7 @@ gfx90a [0x989f];3;8;5;8;30000;0;15;8;1;1;1;1;-2;0;0;0;1;0;0;0;0;0
 gfx90a [0x989f];3;8;6;6;30000;0;12;8;1;8;1;-2;-2;0;0;0;1;0;2;0;0;0
 gfx90a [0x989f];3;8;6;8;30000;0;15;8;1;1;1;-2;-1;0;0;0;1;0;0;2;0;0
 gfx90a [0x989f];3;8;7;7;30000;0;13;8;1;5;1;1;-2;0;0;0;1;0;0;2;0;0
-gfx90a [0x989f];3;8;7;8;30000;0;20;8;1;4;1;-2;0;0;0;0;1;1;0;2;1;0
+gfx90a [0x989f];3;8;7;8;30000;0;20;8;1;6;1;-1;-2;0;0;0;1;0;0;0;0;0
 gfx90a [0x989f];3;8;8;3;30000;0;15;8;1;1;1;1;-1;0;0;0;1;0;2;0;0;0
 gfx90a [0x989f];3;8;8;4;30000;0;15;8;1;7;1;0;-1;0;0;0;1;0;0;2;0;0
 gfx90a [0x989f];3;8;8;5;30000;0;15;8;1;8;1;0;-1;0;0;0;1;0;2;2;0;0
@@ -212,7 +212,7 @@ gfx90a [0x989f];3;9;9;23;30000;0;15;9;1;1;1;-1;-1;0;1;1;1;1;2;2;0;0
 gfx90a [0x989f];3;9;16;9;30000;0;24;9;1;8;1;0;-2;0;0;0;1;0;2;2;0;0
 gfx90a [0x989f];3;9;16;16;30000;0;12;9;1;1;1;1;-2;0;0;0;1;0;2;2;0;0
 gfx90a [0x989f];3;9;23;9;30000;0;26;9;1;8;1;0;-2;0;0;0;1;0;0;2;0;0
-gfx90a [0x989f];3;9;23;23;30000;0;3;8;1;8;1;1;3;0;1;0;1;0;1;2;1;0
+gfx90a [0x989f];3;9;23;23;30000;0;3;8;1;5;1;0;2;0;0;0;1;0;1;2;0;0
 gfx90a [0x989f];3;10;3;3;30000;0;10;10;1;8;1;0;-2;0;1;0;1;1;2;0;0;0
 gfx90a [0x989f];3;10;3;10;30000;0;12;10;1;1;1;-1;-2;0;0;1;1;1;2;0;0;0
 gfx90a [0x989f];3;10;4;4;30000;0;13;10;1;3;1;-2;-1;1;1;0;1;1;0;2;0;0
@@ -249,7 +249,7 @@ gfx90a [0x989f];3;11;11;3;30000;0;22;11;1;2;1;-1;1;0;1;0;1;1;2;0;1;0
 gfx90a [0x989f];3;11;11;4;30000;0;20;11;1;2;1;-1;-2;0;1;1;1;1;0;0;0;0
 gfx90a [0x989f];3;11;11;5;30000;0;20;11;1;5;1;-2;1;0;1;1;1;1;2;0;1;0
 gfx90a [0x989f];3;11;11;6;30000;0;24;11;1;1;1;-2;0;0;0;0;1;1;2;2;0;0
-gfx90a [0x989f];3;11;11;7;30000;0;9;8;1;7;1;0;1;0;1;0;1;0;1;2;0;0
+gfx90a [0x989f];3;11;11;7;30000;0;9;8;1;9;1;1;1;0;0;0;1;0;1;2;0;0
 gfx90a [0x989f];3;11;11;11;30000;0;12;11;1;10;1;-2;-2;0;1;0;1;1;2;0;0;0
 gfx90a [0x989f];3;11;11;16;30000;0;3;8;1;11;1;1;-1;1;0;0;1;0;1;2;0;0
 gfx90a [0x989f];3;11;11;23;30000;0;15;11;1;1;1;1;-1;0;1;1;1;1;2;2;0;0
@@ -351,7 +351,7 @@ gfx90a [0x989f];3;18;23;18;30000;0;4;18;1;4;1;1;3;0;0;0;1;0;1;0;0;0
 gfx90a [0x989f];3;18;23;23;30000;0;4;8;1;16;1;0;3;0;0;0;1;0;1;2;0;0
 gfx90a [0x989f];3;19;19;19;30000;0;40;8;1;10;1;1;3;0;0;0;1;0;1;2;0;0
 gfx90a [0x989f];3;19;19;23;30000;0;40;8;1;15;1;-1;-2;0;0;0;1;0;1;0;0;0
-gfx90a [0x989f];3;23;23;23;30000;0;4;8;1;22;23;-1;3;0;0;0;1;0;1;0;0;0
+gfx90a [0x989f];3;23;23;23;30000;0;4;8;1;22;23;1;0;0;1;0;1;0;1;0;0;0
 gfx90a [0x989f];3;28;28;28;30000;0;3;28;1;28;28;-2;2;0;0;0;1;0;1;0;0;0
 gfx90a [0x989f];3;32;32;32;30000;0;25;32;1;20;1;-2;0;0;1;0;1;0;2;0;0;0
 gfx90a [0x989f];3;35;17;17;30000;0;15;35;1;29;1;1;0;0;1;0;1;0;2;1;0;0
diff --git a/src/acc/opencl/smm/params/tune_multiply_PVC.csv b/src/acc/opencl/smm/params/tune_multiply_PVC.csv
index 5b5a9648737..fe922a84228 100644
--- a/src/acc/opencl/smm/params/tune_multiply_PVC.csv
+++ b/src/acc/opencl/smm/params/tune_multiply_PVC.csv
@@ -1,8 +1,8 @@
 DEVICE;TYPEID;M;N;K;S;GFLOPS;BS;BM;BN;BK;WS;WG;LU;NZ;AL;TB;TC;AP;AA;AB;AC
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;2;2;2;30000;0;8;2;1;1;1;-1;0;0;1;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;3;3;1;30000;0;8;3;1;1;1;1;1;0;0;0;1;0;0;2;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;3;3;3;30000;0;8;3;1;1;1;-1;-2;0;0;0;1;0;2;0;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;4;4;4;30000;0;9;4;1;3;1;-1;1;0;1;0;1;0;2;0;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;3;3;3;30000;0;8;3;1;1;1;-2;-2;0;1;0;1;0;0;0;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;4;4;4;30000;0;8;4;1;3;1;-1;1;0;1;0;1;0;2;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;4;4;5;30000;0;10;4;1;3;1;-1;-2;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;4;4;7;30000;0;9;4;1;3;1;1;0;0;0;0;1;0;2;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;4;4;9;30000;0;8;4;1;2;1;-2;-2;0;0;0;1;0;2;2;0;0
@@ -131,7 +131,7 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;4;28;30000;0;8;5;1;1;5;-2;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;4;32;30000;0;8;5;1;1;1;-2;1;0;0;0;1;0;0;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;4;45;30000;0;8;5;1;1;5;-2;-2;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;5;4;30000;0;12;5;1;1;1;1;-2;0;0;0;1;0;0;2;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;5;5;30000;0;10;5;1;4;1;-1;0;0;1;0;1;0;0;2;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;5;5;30000;0;8;5;1;1;1;1;-1;0;1;0;1;0;2;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;5;7;30000;0;8;5;1;1;1;-2;0;0;0;0;1;0;2;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;5;9;30000;0;8;5;1;1;1;-1;-2;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;5;13;30000;0;8;5;1;1;1;-2;1;0;0;0;1;0;0;0;0;0
@@ -143,7 +143,7 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;5;26;30000;0;8;5;1;1;1;-2;-2;0;0;
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;5;28;30000;0;8;5;1;1;5;-2;-2;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;5;32;30000;0;8;5;1;1;1;-2;-1;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;6;20;30000;0;8;5;1;1;6;-1;-1;0;0;0;1;0;0;0;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;6;32;30000;0;8;5;1;1;6;1;1;0;0;0;1;0;0;0;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;6;32;30000;0;8;5;1;1;6;1;1;0;1;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;7;4;30000;0;12;5;1;1;1;-2;1;0;0;0;1;0;2;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;7;5;30000;0;12;5;1;1;1;-1;0;0;0;0;1;0;2;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;7;7;30000;0;10;5;1;1;1;-2;-1;0;0;0;1;0;0;2;0;0
@@ -196,7 +196,7 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;25;26;30000;0;8;5;1;1;25;-2;-2;0;
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;25;28;30000;0;8;5;1;1;25;-1;0;0;0;0;1;1;2;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;25;32;30000;0;8;5;1;1;25;1;0;0;0;0;1;0;0;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;25;45;30000;0;8;5;1;1;25;-1;0;0;0;0;1;1;0;0;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;26;4;30000;0;18;5;1;1;26;-2;1;0;0;0;1;0;0;0;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;26;4;30000;0;18;5;1;1;26;-2;0;0;0;0;1;0;0;1;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;26;5;30000;0;18;5;1;1;1;-2;-2;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;26;7;30000;0;14;5;1;1;26;1;-2;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;26;9;30000;0;14;5;1;1;26;1;0;0;0;0;1;0;2;2;0;0
@@ -219,9 +219,9 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;28;26;30000;0;8;5;1;1;28;0;-2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;28;28;30000;0;8;5;1;1;28;-2;0;1;1;1;1;1;2;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;28;32;30000;0;8;5;1;1;28;-2;-2;1;1;1;1;1;2;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;28;45;30000;0;8;5;1;1;28;1;-2;0;1;0;1;1;2;0;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;32;4;30000;0;18;5;1;5;1;0;-1;0;0;0;1;0;0;2;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;32;4;30000;0;18;5;1;1;1;-2;1;0;0;0;1;0;2;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;32;5;30000;0;17;5;1;1;1;1;-2;0;0;0;1;0;2;1;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;32;6;30000;0;18;5;1;2;32;0;-1;0;1;0;1;0;2;2;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;32;6;30000;0;18;5;1;2;32;-1;-2;0;1;0;1;0;2;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;32;7;30000;0;15;5;1;1;32;-2;-2;0;0;0;1;0;0;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;32;9;30000;0;15;5;1;1;32;1;-1;0;0;0;1;0;2;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;32;13;30000;0;15;5;1;1;1;1;0;0;0;0;1;0;2;2;0;0
@@ -243,9 +243,9 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;45;26;30000;0;6;5;1;1;45;-2;-1;0;
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;45;28;30000;0;5;5;1;1;45;-1;-2;1;0;1;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;45;32;30000;0;4;5;1;1;45;-1;-2;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;5;45;45;30000;0;9;5;1;1;45;-2;-1;0;0;0;1;1;0;2;0;1
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;5;20;30000;0;8;6;1;1;6;-2;-1;0;0;0;1;0;0;0;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;5;20;30000;0;8;6;1;1;6;-2;-1;0;1;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;5;32;30000;0;8;6;1;1;6;-2;-1;0;0;0;1;0;0;0;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;6;6;30000;0;9;6;1;3;1;-2;0;0;1;0;1;1;0;0;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;6;6;30000;0;9;6;1;4;1;-2;0;0;1;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;6;7;30000;0;10;6;1;1;1;-1;-1;0;0;0;1;0;0;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;6;8;30000;0;10;6;1;1;1;-2;-1;0;0;0;1;0;0;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;6;20;30000;0;10;6;1;1;6;-2;-2;0;1;0;1;0;0;0;0;0
@@ -258,12 +258,12 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;8;7;30000;0;14;6;1;1;1;-1;-2;0;0;
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;8;8;30000;0;13;6;1;1;1;-1;0;0;0;0;1;0;2;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;13;20;30000;0;10;6;1;1;13;1;1;0;0;0;1;0;0;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;13;32;30000;0;8;6;1;1;13;1;1;0;0;0;1;0;2;0;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;20;5;30000;0;18;6;1;1;20;-1;1;0;0;0;1;0;0;1;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;20;5;30000;0;17;6;1;1;20;-1;1;0;0;0;1;0;0;1;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;20;6;30000;0;15;6;1;1;20;-2;1;0;0;0;1;0;0;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;20;13;30000;0;14;6;1;1;20;-1;-1;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;20;20;30000;0;8;6;1;1;20;-1;-1;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;20;32;30000;0;8;6;1;1;20;0;-2;0;0;0;1;0;0;0;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;32;5;30000;0;20;6;1;6;32;1;-1;0;0;0;1;0;0;0;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;32;5;30000;0;22;6;1;6;32;1;-1;0;1;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;32;6;30000;0;18;6;1;6;32;1;1;0;0;0;1;0;2;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;32;13;30000;0;15;6;1;1;32;1;1;0;0;0;1;0;2;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;6;32;20;30000;0;8;6;1;1;32;-1;1;0;0;0;1;0;0;0;0;0
@@ -294,7 +294,7 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;7;6;8;30000;0;10;7;1;1;1;1;1;0;0;0;
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;7;7;4;30000;0;14;7;1;1;1;1;-1;0;0;0;1;0;0;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;7;7;5;30000;0;16;7;1;1;1;-1;0;0;0;0;1;0;0;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;7;7;6;30000;0;16;7;1;1;1;1;-1;0;0;0;1;0;0;0;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;7;7;7;30000;0;9;7;1;3;1;0;-2;0;1;0;1;0;0;2;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;7;7;7;30000;0;9;7;1;3;1;0;-1;0;1;0;1;0;2;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;7;7;8;30000;0;14;7;1;1;1;1;-1;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;7;7;9;30000;0;13;7;1;1;1;-2;-1;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;7;7;13;30000;0;10;7;1;1;1;-1;-2;0;0;0;1;0;0;2;0;0
@@ -358,7 +358,7 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;4;45;30000;0;8;9;1;1;9;-1;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;5;4;30000;0;16;9;1;1;1;-2;1;0;0;0;1;0;0;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;5;5;30000;0;16;9;1;1;1;0;1;0;0;0;1;0;2;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;5;7;30000;0;16;9;1;1;1;-2;-2;1;0;1;1;0;0;2;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;5;9;30000;0;11;9;1;1;1;-1;-1;1;0;1;1;0;0;2;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;5;9;30000;0;11;9;1;1;1;-1;-2;0;0;0;1;0;0;1;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;5;13;30000;0;10;9;1;1;9;-1;1;0;0;0;1;0;2;1;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;5;25;30000;0;8;9;1;1;9;-2;0;0;0;0;1;0;0;1;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;5;26;30000;0;8;9;1;3;9;-1;1;0;0;0;1;1;0;2;1;0
@@ -432,8 +432,8 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;28;13;30000;0;15;9;1;2;28;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;28;25;30000;0;8;9;1;1;28;0;0;0;1;0;1;0;1;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;28;26;30000;0;10;9;1;1;28;-2;0;0;1;1;1;0;0;0;1;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;28;28;30000;0;12;9;1;1;28;0;0;0;0;0;1;0;2;0;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;28;32;30000;0;15;9;1;4;28;-2;-1;0;0;0;1;0;1;0;0;1
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;28;45;30000;0;15;9;1;1;28;1;0;0;0;0;1;0;0;0;0;1
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;28;32;30000;0;15;9;1;8;28;-2;0;0;0;0;1;0;1;0;0;1
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;28;45;30000;0;15;9;1;1;28;-1;0;0;0;0;1;0;0;0;0;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;32;4;30000;0;39;9;1;1;32;-1;-1;0;0;0;1;0;2;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;32;5;30000;0;38;9;1;1;32;-1;-1;0;0;0;1;0;2;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;9;32;7;30000;0;25;9;1;8;32;0;0;0;0;0;1;0;1;2;0;0
@@ -517,7 +517,7 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;13;24;30000;0;12;13;1;1;13;-1;-1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;13;25;30000;0;8;13;1;1;13;-1;0;1;0;1;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;13;26;30000;0;14;13;1;1;13;-1;-2;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;13;28;30000;0;8;13;1;1;13;-1;-1;1;1;0;1;1;0;0;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;13;32;30000;0;8;13;1;1;1;-2;1;0;0;0;1;0;0;0;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;13;32;30000;0;8;13;1;1;1;-2;0;0;0;0;1;0;0;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;13;45;30000;0;8;13;1;1;13;-2;0;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;14;13;30000;0;15;13;1;1;14;-2;1;0;0;0;1;0;0;0;0;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;14;14;30000;0;15;13;1;1;14;-2;1;0;0;0;1;0;2;0;0;0
@@ -533,7 +533,7 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;17;32;30000;0;8;13;1;1;1;1;-2;1;
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;20;5;30000;0;17;13;1;8;20;0;1;0;0;0;1;0;1;1;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;20;6;30000;0;9;13;1;12;20;-1;-2;0;0;0;1;0;1;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;20;13;30000;0;11;13;1;9;20;1;-1;0;0;0;1;0;1;0;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;20;20;30000;0;17;8;1;6;20;0;-1;0;0;0;1;0;1;2;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;20;20;30000;0;15;8;1;6;20;1;-1;0;0;0;1;0;1;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;20;32;30000;0;15;8;1;1;20;-2;1;0;0;0;1;0;2;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;24;5;30000;0;22;13;1;5;1;1;-1;1;1;0;1;0;1;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;24;13;30000;0;13;13;1;10;24;-1;0;0;0;0;1;0;1;0;0;0
@@ -559,23 +559,23 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;26;13;30000;0;11;13;1;13;26;-2;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;26;14;30000;0;11;13;1;8;26;-2;-1;0;0;0;1;0;1;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;26;24;30000;0;15;13;1;1;26;1;1;0;0;0;1;0;0;0;0;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;26;25;30000;0;15;13;1;1;26;-1;0;1;0;0;1;0;0;0;0;1
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;26;26;30000;0;15;13;1;1;26;1;1;0;0;0;1;0;0;0;0;1
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;26;26;30000;0;15;13;1;1;26;1;-2;0;0;0;1;0;0;0;0;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;26;28;30000;0;15;13;1;1;26;0;1;0;0;0;1;0;0;0;0;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;26;32;30000;0;15;13;1;1;26;-1;-1;0;0;0;1;0;0;0;0;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;26;45;30000;0;15;13;1;1;26;-1;1;0;0;0;1;0;0;0;0;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;28;4;30000;0;25;13;1;4;28;-1;-1;0;0;0;1;0;2;1;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;28;5;30000;0;17;13;1;9;28;0;0;0;0;0;1;0;1;1;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;28;5;30000;0;17;13;1;11;28;0;0;0;1;0;1;0;0;1;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;28;7;30000;0;21;13;1;11;28;1;0;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;28;9;30000;0;16;13;1;10;28;0;0;0;0;0;1;0;2;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;28;13;30000;0;13;13;1;11;28;1;1;0;0;0;1;0;1;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;28;25;30000;0;15;13;1;1;28;-1;-2;0;0;0;1;0;0;0;0;1
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;28;26;30000;0;15;13;1;4;28;1;-2;0;0;0;1;0;1;0;0;1
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;28;28;30000;0;15;13;1;1;28;1;0;0;0;0;1;0;0;0;0;1
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;28;26;30000;0;15;13;1;5;28;1;-2;0;1;0;1;0;1;0;0;1
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;28;28;30000;0;15;13;1;1;28;1;1;0;0;0;1;0;0;0;0;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;28;32;30000;0;15;13;1;1;28;1;-2;0;0;0;1;0;0;0;0;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;28;45;30000;0;15;13;1;1;28;0;1;0;0;0;1;0;0;0;0;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;32;4;30000;0;25;13;1;7;1;-1;0;0;0;0;1;0;1;1;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;32;5;30000;0;25;13;1;11;1;0;-2;0;0;0;1;0;1;1;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;32;6;30000;0;25;13;1;5;32;-2;0;0;0;0;1;0;2;2;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;32;6;30000;0;25;13;1;9;32;-2;0;0;0;0;1;0;2;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;32;7;30000;0;25;13;1;4;32;0;0;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;32;9;30000;0;16;13;1;7;32;-2;0;0;0;0;1;0;2;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;32;13;30000;0;13;13;1;13;1;-1;-1;0;0;0;1;0;1;2;0;0
@@ -592,7 +592,7 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;45;4;30000;0;30;13;1;11;45;1;1;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;45;5;30000;0;30;8;1;6;45;-2;0;0;0;0;1;0;1;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;45;7;30000;0;37;13;1;6;45;1;0;0;0;0;1;0;2;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;45;9;30000;0;29;13;1;13;45;1;-2;0;0;0;1;0;2;0;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;45;13;30000;0;15;13;1;3;45;-2;0;0;0;0;1;0;1;0;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;45;13;30000;0;15;13;1;6;45;-1;0;0;0;0;1;0;1;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;45;25;30000;0;17;8;1;5;45;-2;2;0;0;0;1;0;1;2;0;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;45;26;30000;0;30;13;1;1;45;-1;-2;1;0;0;1;0;2;2;0;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;13;45;28;30000;0;10;8;1;6;45;-1;-1;0;1;0;1;0;1;2;1;1
@@ -738,7 +738,7 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;24;5;5;30000;0;17;8;1;21;1;-1;0;1;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;24;5;13;30000;0;18;8;1;4;1;-1;0;0;0;0;1;0;2;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;24;5;16;30000;0;8;8;1;7;1;-1;4;1;1;0;1;0;2;2;1;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;24;5;24;30000;0;8;8;1;1;24;0;-2;1;1;0;1;1;0;0;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;24;5;26;30000;0;9;8;1;1;24;-1;3;0;0;0;1;0;0;0;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;24;5;26;30000;0;8;8;1;1;24;-1;3;0;1;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;24;5;32;30000;0;8;8;1;1;24;0;-1;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;24;13;5;30000;0;20;24;1;2;1;-2;-1;0;0;0;1;0;2;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;24;13;13;30000;0;15;8;1;13;24;-1;0;0;0;0;1;0;1;2;0;0
@@ -1039,7 +1039,7 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;28;13;25;30000;0;17;16;1;22;28;-2;2
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;28;13;26;30000;0;17;16;1;19;28;0;2;0;0;0;1;0;2;1;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;28;13;28;30000;0;15;8;1;1;28;-2;1;0;0;0;1;0;2;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;28;13;32;30000;0;15;8;1;1;28;1;-2;0;0;0;1;0;2;0;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;28;13;45;30000;0;30;8;1;14;28;1;-1;0;0;0;1;0;2;2;0;1
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;28;13;45;30000;0;3;8;1;12;28;1;4;0;0;0;1;0;2;1;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;28;25;4;30000;0;30;8;1;22;28;-1;4;0;0;0;1;0;1;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;28;25;5;30000;0;30;8;1;21;28;-1;-2;0;0;0;1;0;1;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;28;25;7;30000;0;30;8;1;5;28;0;4;0;0;0;1;0;1;2;0;0
@@ -1110,7 +1110,7 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;29;32;32;30000;0;30;29;1;1;32;-2;5;
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;29;32;55;30000;0;5;16;1;1;32;1;3;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;29;55;16;30000;0;15;29;1;12;55;0;1;0;0;0;1;0;1;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;29;55;29;30000;0;27;8;1;1;55;-2;4;0;0;0;1;0;2;0;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;29;55;32;30000;0;14;16;1;24;55;0;3;0;0;0;1;0;1;1;0;1
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;29;55;32;30000;0;6;16;1;24;55;0;-2;0;0;0;1;0;1;1;0;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;29;55;55;30000;0;3;16;1;1;55;0;2;0;0;0;1;0;2;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;30;30;30;30000;0;8;30;1;1;1;-2;1;0;0;1;1;1;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;30;30;76;30000;0;8;30;1;1;30;-1;1;0;0;0;1;0;0;0;0;0
@@ -1184,11 +1184,11 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;17;5;30000;0;29;8;1;31;1;1;1;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;17;13;30000;0;27;8;1;23;1;0;-1;0;0;0;1;0;1;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;17;17;30000;0;15;32;1;1;1;-1;1;1;0;1;1;1;2;0;0;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;17;32;30000;0;7;8;1;1;1;-2;4;0;0;0;1;0;2;1;0;1
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;20;5;30000;0;40;8;1;29;32;-2;4;0;0;0;1;0;0;0;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;20;5;30000;0;40;8;1;30;32;-2;4;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;20;6;30000;0;31;32;1;4;32;-2;4;0;0;0;1;0;0;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;20;13;30000;0;27;32;1;2;32;-1;4;0;0;0;1;0;1;2;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;20;20;30000;0;12;32;1;1;32;1;4;0;0;0;1;0;0;0;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;20;32;30000;0;3;32;1;22;32;-2;4;0;0;0;1;0;0;0;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;20;20;30000;0;12;8;1;1;32;1;-1;0;0;0;1;0;0;0;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;20;32;30000;0;2;8;1;1;32;1;4;0;1;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;22;9;30000;0;27;8;1;8;32;1;-1;0;0;0;1;0;1;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;22;22;30000;0;12;8;1;1;1;1;4;1;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;22;32;30000;0;9;8;1;1;1;-1;-1;0;0;0;1;0;1;1;0;1
@@ -1253,7 +1253,7 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;32;32;30000;0;8;32;1;1;1;-1;1;0;
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;32;45;30000;0;10;32;1;1;32;-1;1;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;32;55;30000;0;10;32;1;1;32;-1;1;0;0;0;1;0;0;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;45;4;30000;0;61;8;1;6;45;0;-1;0;0;0;1;0;0;2;0;0
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;45;5;30000;0;30;8;1;22;45;-2;2;0;0;0;1;0;1;2;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;45;5;30000;0;30;16;1;20;45;-2;2;0;0;0;1;0;1;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;45;7;30000;0;28;32;1;2;45;-1;4;0;0;0;1;0;2;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;45;9;30000;0;27;16;1;23;45;-1;-1;0;0;0;1;0;1;2;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;32;45;13;30000;0;28;8;1;5;45;-2;-1;0;0;0;1;0;1;2;0;0
@@ -1272,7 +1272,7 @@ Intel(R) Data Center GPU Max 1550 [0x0bd5];3;35;35;32;30000;0;9;8;1;21;35;0;0;0;
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;35;35;35;30000;0;9;8;1;34;1;0;4;0;1;1;1;1;2;0;1;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;36;36;36;30000;0;6;8;1;23;36;1;5;1;0;0;1;0;2;2;1;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;40;40;40;30000;0;12;8;1;1;1;0;0;1;0;0;1;1;1;0;0;1
-Intel(R) Data Center GPU Max 1550 [0x0bd5];3;45;4;4;30000;0;3;45;1;36;45;0;2;0;1;0;1;0;0;0;0;0
+Intel(R) Data Center GPU Max 1550 [0x0bd5];3;45;4;4;30000;0;6;24;1;17;45;-2;5;0;1;0;1;0;2;0;0;0
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;45;13;13;30000;0;17;8;1;4;45;1;1;0;0;0;1;0;2;2;0;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;45;32;28;30000;0;16;8;1;7;45;-2;3;0;0;0;1;0;1;0;0;1
 Intel(R) Data Center GPU Max 1550 [0x0bd5];3;45;45;45;30000;0;30;45;1;1;45;-2;1;0;0;0;1;0;2;0;0;1