From 5b0a12e3c9c82597a92bd8b43b430c95cdb0ca7a Mon Sep 17 00:00:00 2001 From: Hans Pabst Date: Sun, 10 Mar 2024 18:17:34 +0100 Subject: [PATCH] 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:". --- src/acc/acc_libsmm.h | 2 +- src/acc/cuda_hip/calculate_norms.cpp | 3 ++- src/acc/opencl/acc_opencl.c | 14 +++++++++++--- src/acc/opencl/acc_opencl.h | 14 +++++++++----- src/acc/opencl/smm/opencl_libsmm.c | 2 +- 5 files changed, 24 insertions(+), 11 deletions(-) diff --git a/src/acc/acc_libsmm.h b/src/acc/acc_libsmm.h index ddf3c1ccdba..06957d74074 100644 --- 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 index c827bb7b190..fc69a80dcd8 100644 --- 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 index b8f67f50c91..64756cc7d11 100644 --- 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 index c109de21ef3..8cdbc15f2df 100644 --- 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 index e1384dedf73..6089fe572bd 100644 --- 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);