Fixed interface for calculate_norms

* Calculate_norms: const-correctness to match Fortran wrapper's INTENT. * OpenCL BE: Made ACC_OPENCL_WA flag-based (bit-wise toggles). * OpenCL BE: Fixed macro order (acc_opencl.h). * Eventually suppress "Fortran runtime warning:".
cp2k · Mar 10, 2024 · 5b0a12e · 5b0a12e
1 parent ab07d39
commit 5b0a12e
Show file tree

Hide file tree

Showing 5 changed files with 24 additions and 11 deletions.
diff --git a/src/acc/acc_libsmm.h b/src/acc/acc_libsmm.h
@@ -46,7 +46,7 @@ int libsmm_acc_process(const int* host_param_stack, const int* dev_param_stack,
   const void* dev_a_data, const void* dev_b_data, void* dev_c_data, int m_max, int n_max, int k_max, int max_kernel_dim,
   c_dbcsr_acc_bool_t def_mnk, void* stack_stream, void* c_stream);
 
-int c_calculate_norms(double* mat, int nblks, int* offsets, int* nelems, float* norms, void* stream_ptr);
+int c_calculate_norms(const double* mat, int nblks, const int* offsets, const int* nelems, float* norms, void* stream_ptr);
 
 static const char libsmm_acc_transpose_routine_name_str[] = "jit_kernel_transpose";
 static const char* const libsmm_acc_transpose_routine_name_ptr = libsmm_acc_transpose_routine_name_str;

diff --git a/src/acc/cuda_hip/calculate_norms.cpp b/src/acc/cuda_hip/calculate_norms.cpp
@@ -95,7 +95,8 @@ __global__ void calculate_norms_d(
   }
 }
 
-extern "C" int c_calculate_norms(double* mat, int nblks, int* offsets, int* nelems, float* norms, void* stream_ptr) {
+extern "C" int c_calculate_norms(
+  const double* mat, int nblks, const int* offsets, const int* nelems, float* norms, void* stream_ptr) {
   int warp_size = acc_get_gpu_warp_size();
 
   dim3 grid(nblks);

diff --git a/src/acc/opencl/acc_opencl.c b/src/acc/opencl/acc_opencl.c
@@ -67,6 +67,14 @@ c_dbcsr_acc_opencl_config_t c_dbcsr_acc_opencl_config;
 int c_dbcsr_acc_opencl_active_id;
 #  endif
 
+
+void __wrap__gfortran_runtime_warning_at(const char* /*where*/, const char* /*message*/, ...);
+void __wrap__gfortran_runtime_warning_at(const char* where, const char* message, ...) {
+  LIBXSMM_UNUSED(message);
+  LIBXSMM_UNUSED(where);
+}
+
+
 void c_dbcsr_acc_opencl_notify(const char /*errinfo*/[], const void* /*private_info*/, size_t /*cb*/, void* /*user_data*/);
 void c_dbcsr_acc_opencl_notify(const char errinfo[], const void* private_info, size_t cb, void* user_data) {
   LIBXSMM_UNUSED(private_info);
@@ -215,7 +223,7 @@ int c_dbcsr_acc_init(void) {
     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)), wa = neo * (NULL == env_wa ? 2 : atoi(env_wa));
+    const int neo = (NULL == env_neo ? 1 : atoi(env_neo)), wa = neo * (NULL == env_wa ? 3 : atoi(env_wa));
 #  if defined(ACC_OPENCL_ASYNC)
     const char* const env_async = (ACC_OPENCL_ASYNC);
     const int async_default = 3;
@@ -296,10 +304,10 @@ int c_dbcsr_acc_init(void) {
       static char* key_value[] = {
         "NEOReadDebugKeys=1", "EnableRecoverablePageFaults=0", "DirectSubmissionOverrideBlitterSupport=0"};
       if (NULL == env_neo) ACC_OPENCL_EXPECT(0 == LIBXSMM_PUTENV(key_value[0]));
-      if (NULL == getenv("EnableRecoverablePageFaults")) {
+      if (0 != (1 & wa) && NULL == getenv("EnableRecoverablePageFaults")) {
         ACC_OPENCL_EXPECT(0 == LIBXSMM_PUTENV(key_value[1]));
       }
-      if (NULL == getenv("DirectSubmissionOverrideBlitterSupport") && 2 <= wa) {
+      if (0 != (2 & wa) && NULL == getenv("DirectSubmissionOverrideBlitterSupport")) {
         ACC_OPENCL_EXPECT(0 == LIBXSMM_PUTENV(key_value[2]));
       }
     }

diff --git a/src/acc/opencl/acc_opencl.h b/src/acc/opencl/acc_opencl.h
@@ -9,6 +9,11 @@
 #ifndef ACC_OPENCL_H
 #define ACC_OPENCL_H
 
+/* Support for other libraries, e.g., CP2K's DBM/DBT */
+#if defined(__OFFLOAD_OPENCL) && !defined(__OPENCL)
+#  define __OPENCL
+#endif
+
 #if defined(__OPENCL)
 #  if !defined(CL_TARGET_OPENCL_VERSION)
 #    define CL_TARGET_OPENCL_VERSION 220
@@ -116,6 +121,10 @@
 #  define ACC_OPENCL_PROFILE
 #endif
 
+#if defined(__OFFLOAD_OPENCL) && !defined(ACC_OPENCL_MEM_DEVPTR)
+#  error Support for ACC_OPENCL_MEM_DEVPTR is required!
+#endif
+
 /* attaching c_dbcsr_acc_opencl_stream_t is needed */
 #define ACC_OPENCL_STREAM(A) ((const c_dbcsr_acc_opencl_stream_t*)(A))
 /* incompatible with c_dbcsr_acc_event_record */
@@ -161,11 +170,6 @@
     clCreateCommandQueue(CTX, DEV, (cl_command_queue_properties)(NULL != (PROPS) ? ((PROPS)[1]) : 0), RESULT)
 #endif
 
-/* Support for other libraries, e.g., CP2K's DBM/DBT */
-#if defined(ACC_OPENCL_MEM_DEVPTR) && defined(__OFFLOAD_OPENCL) && !defined(__OPENCL)
-#  define __OPENCL
-#endif
-
 #if LIBXSMM_VERSION4(1, 17, 0, 0) < LIBXSMM_VERSION_NUMBER
 #  define ACC_OPENCL_EXPECT(EXPR) LIBXSMM_EXPECT(EXPR)
 #  define LIBXSMM_STRISTR libxsmm_stristr

diff --git a/src/acc/opencl/smm/opencl_libsmm.c b/src/acc/opencl/smm/opencl_libsmm.c
@@ -1563,7 +1563,7 @@ int libsmm_acc_process(const int* host_param_stack, const int* dev_param_stack,
 }
 
 
-int c_calculate_norms(double* mat, int nblks, int* offsets, int* nelems, float* norms, void* stream_ptr) {
+int c_calculate_norms(const double* mat, int nblks, const int* offsets, const int* nelems, float* norms, void* stream_ptr) {
   LIBXSMM_UNUSED(mat);
   LIBXSMM_UNUSED(nblks);
   LIBXSMM_UNUSED(offsets);