From cc5f6914f6e92688b59ee187560a5d6a3a5ace48 Mon Sep 17 00:00:00 2001 From: Anthony Scemama Date: Wed, 6 Apr 2022 16:26:35 +0200 Subject: [PATCH] Cleaning --- org/qmckl_jastrow.org | 1248 +++++++++++++++++++++-------------------- org/qmckl_mo.org | 2 - 2 files changed, 646 insertions(+), 604 deletions(-) diff --git a/org/qmckl_jastrow.org b/org/qmckl_jastrow.org index 35003f5..666da47 100644 --- a/org/qmckl_jastrow.org +++ b/org/qmckl_jastrow.org @@ -151,7 +151,6 @@ int main() { | ~factor_en_deriv_e_date~ | ~uint64_t~ | out | Keep track of the date for the en derivative | | ~factor_een_deriv_e~ | ~double[4][nelec][walk_num]~ | out | Derivative of the Jastrow factor: electron-electron-nucleus part | | ~factor_een_deriv_e_date~ | ~uint64_t~ | out | Keep track of the date for the een derivative | - | ~offload_type~ | ~qmckl_jastrow_offload_type~ | in | Enum type to change offload type at runtime | computed data: @@ -328,14 +327,6 @@ kappa_inv = 1.0/kappa ** Data structure -#+begin_src c :comments org :tangle (eval h_type) -typedef enum qmckl_jastrow_offload_type{ - OFFLOAD_NONE, - OFFLOAD_OPENACC, - OFFLOAD_CUBLAS -} qmckl_jastrow_offload_type; -#+end_src - #+begin_src c :comments org :tangle (eval h_private_type) typedef struct qmckl_jastrow_struct{ int32_t uninitialized; @@ -381,7 +372,10 @@ typedef struct qmckl_jastrow_struct{ uint64_t een_rescaled_n_deriv_e_date; bool provided; char * type; - qmckl_jastrow_offload_type offload_type; + + #ifdef HAVE_HPC + bool gpu_offload; + #endif } qmckl_jastrow_struct; #+end_src @@ -426,7 +420,6 @@ qmckl_exit_code qmckl_get_jastrow_type_nucl_vector (qmckl_context context, int qmckl_exit_code qmckl_get_jastrow_aord_vector (qmckl_context context, double * const aord_vector, const int64_t size_max); qmckl_exit_code qmckl_get_jastrow_bord_vector (qmckl_context context, double * const bord_vector, const int64_t size_max); qmckl_exit_code qmckl_get_jastrow_cord_vector (qmckl_context context, double * const cord_vector, const int64_t size_max); -qmckl_exit_code qmckl_get_jastrow_offload_type (qmckl_context context, qmckl_jastrow_offload_type * const offload_type); #+end_src Along with these core functions, calculation of the jastrow factor @@ -724,32 +717,6 @@ qmckl_get_jastrow_cord_vector (const qmckl_context context, return QMCKL_SUCCESS; } -qmckl_exit_code qmckl_get_jastrow_offload_type (const qmckl_context context, qmckl_jastrow_offload_type* const offload_type) { - - if (qmckl_context_check(context) == QMCKL_NULL_CONTEXT) { - return (char) 0; - } - - if (offload_type == NULL) { - return qmckl_failwith( context, - QMCKL_INVALID_ARG_2, - "qmckl_get_jastrow_offload_type", - "offload_type is a null pointer"); - } - - qmckl_context_struct* const ctx = (qmckl_context_struct*) context; - assert (ctx != NULL); - - int32_t mask = 1 << 0; - - if ( (ctx->jastrow.uninitialized & mask) != 0) { - return QMCKL_NOT_PROVIDED; - } - - *offload_type = ctx->jastrow.offload_type; - return QMCKL_SUCCESS; -} - #+end_src ** Initialization functions @@ -764,7 +731,6 @@ qmckl_exit_code qmckl_set_jastrow_type_nucl_vector (qmckl_context context, con qmckl_exit_code qmckl_set_jastrow_aord_vector (qmckl_context context, const double * aord_vector, const int64_t size_max); qmckl_exit_code qmckl_set_jastrow_bord_vector (qmckl_context context, const double * bord_vector, const int64_t size_max); qmckl_exit_code qmckl_set_jastrow_cord_vector (qmckl_context context, const double * cord_vector, const int64_t size_max); -qmckl_exit_code qmckl_set_jastrow_offload_type (qmckl_context context, const qmckl_jastrow_offload_type offload_type); #+end_src #+NAME:pre2 @@ -1101,14 +1067,6 @@ qmckl_set_jastrow_cord_vector(qmckl_context context, <> } -qmckl_exit_code -qmckl_set_jastrow_offload_type(qmckl_context context, const qmckl_jastrow_offload_type offload_type) -{ -<> - ctx->jastrow.offload_type = offload_type; - return QMCKL_SUCCESS; -} - #+end_src When the required information is completely entered, other data structures are @@ -1155,6 +1113,13 @@ qmckl_exit_code qmckl_finalize_jastrow(qmckl_context context) { NULL); } + /* Decide if the Jastrow if offloaded on GPU or not */ +#if defined(HAVE_HPC) && (defined(HAVE_CUBLAS_OFFLOAD) || defined(HAVE_OPENACC_OFFLOAD) || defined(HAVE_OPENMP_OFFLOAD)) + ctx->jastrow.gpu_offload = true; // ctx->electron.num > 100; +#else + ctx->jastrow.gpu_offload = false; +#endif + qmckl_exit_code rc = QMCKL_SUCCESS; return rc; @@ -1540,16 +1505,16 @@ qmckl_exit_code qmckl_compute_asymp_jasb ( } #+end_src - #+CALL: generate_c_header(table=qmckl_asymp_jasb_args,rettyp=get_value("CRetType"),fname=get_value("Name")) +# #+CALL: generate_c_header(table=qmckl_asymp_jasb_args,rettyp=get_value("CRetType"),fname=get_value("Name")) #+RESULTS: - #+begin_src c :tangle (eval h_func) :comments org + #+begin_src c :comments org :tangle (eval h_private_func) :noweb yes :exports none qmckl_exit_code qmckl_compute_asymp_jasb ( - const qmckl_context context, - const int64_t bord_num, - const double* bord_vector, - const double rescale_factor_kappa_ee, - double* const asymp_jasb ); + const qmckl_context context, + const int64_t bord_num, + const double* bord_vector, + const double rescale_factor_kappa_ee, + double* const asymp_jasb ); #+end_src @@ -1892,19 +1857,19 @@ qmckl_exit_code qmckl_compute_factor_ee ( #+end_src #+CALL: generate_c_header(table=qmckl_factor_ee_args,rettyp=get_value("CRetType"),fname=get_value("Name")) - + #+RESULTS: - #+begin_src c :tangle (eval h_func) :comments org + #+begin_src c :comments org :tangle (eval h_private_func) :noweb yes :exports none qmckl_exit_code qmckl_compute_factor_ee ( - const qmckl_context context, - const int64_t walk_num, - const int64_t elec_num, - const int64_t up_num, - const int64_t bord_num, - const double* bord_vector, - const double* ee_distance_rescaled, - const double* asymp_jasb, - double* const factor_ee ); + const qmckl_context context, + const int64_t walk_num, + const int64_t elec_num, + const int64_t up_num, + const int64_t bord_num, + const double* bord_vector, + const double* ee_distance_rescaled, + const double* asymp_jasb, + double* const factor_ee ); #+end_src @@ -2202,21 +2167,21 @@ integer function qmckl_compute_factor_ee_deriv_e_f( & end function qmckl_compute_factor_ee_deriv_e_f #+end_src - #+CALL: generate_c_header(table=qmckl_factor_ee_deriv_e_args,rettyp=get_value("CRetType"),fname=get_value("Name")) +# #+CALL: generate_c_header(table=qmckl_factor_ee_deriv_e_args,rettyp=get_value("CRetType"),fname=get_value("Name")) #+RESULTS: - #+begin_src c :tangle (eval h_func) :comments org + #+begin_src c :comments org :tangle (eval h_private_func) :noweb yes :exports none qmckl_exit_code qmckl_compute_factor_ee_deriv_e ( - const qmckl_context context, - const int64_t walk_num, - const int64_t elec_num, - const int64_t up_num, - const int64_t bord_num, - const double* bord_vector, - const double* ee_distance_rescaled, - const double* ee_distance_rescaled_deriv_e, - const double* asymp_jasb, - double* const factor_ee_deriv_e ); + const qmckl_context context, + const int64_t walk_num, + const int64_t elec_num, + const int64_t up_num, + const int64_t bord_num, + const double* bord_vector, + const double* ee_distance_rescaled, + const double* ee_distance_rescaled_deriv_e, + const double* asymp_jasb, + double* const factor_ee_deriv_e ); #+end_src @@ -2224,8 +2189,8 @@ end function qmckl_compute_factor_ee_deriv_e_f #+RESULTS: #+begin_src f90 :tangle (eval f) :comments org :exports none -integer(c_int32_t) function qmckl_compute_factor_ee_deriv_e & - (context, & + integer(c_int32_t) function qmckl_compute_factor_ee_deriv_e & + (context, & walk_num, & elec_num, & up_num, & @@ -2235,7 +2200,7 @@ integer(c_int32_t) function qmckl_compute_factor_ee_deriv_e & ee_distance_rescaled_deriv_e, & asymp_jasb, & factor_ee_deriv_e) & - bind(C) result(info) + bind(C) result(info) use, intrinsic :: iso_c_binding implicit none @@ -2245,7 +2210,7 @@ integer(c_int32_t) function qmckl_compute_factor_ee_deriv_e & integer (c_int64_t) , intent(in) , value :: elec_num integer (c_int64_t) , intent(in) , value :: up_num integer (c_int64_t) , intent(in) , value :: bord_num - real (c_double ) , intent(in) :: bord_vector(bord_num + 1) + real (c_double ) , intent(in) :: bord_vector(bord_num+1) real (c_double ) , intent(in) :: ee_distance_rescaled(elec_num,elec_num,walk_num) real (c_double ) , intent(in) :: ee_distance_rescaled_deriv_e(elec_num,elec_num,4,walk_num) real (c_double ) , intent(in) :: asymp_jasb(2) @@ -2253,7 +2218,7 @@ integer(c_int32_t) function qmckl_compute_factor_ee_deriv_e & integer(c_int32_t), external :: qmckl_compute_factor_ee_deriv_e_f info = qmckl_compute_factor_ee_deriv_e_f & - (context, & + (context, & walk_num, & elec_num, & up_num, & @@ -2676,21 +2641,20 @@ qmckl_exit_code qmckl_compute_factor_en ( #+end_src - #+CALL: generate_c_header(table=qmckl_factor_en_args,rettyp=get_value("CRetType"),fname=get_value("Name")) +# #+CALL: generate_c_header(table=qmckl_factor_en_args,rettyp=get_value("CRetType"),fname=get_value("Name")) - #+RESULTS: - #+begin_src c :tangle (eval h_func) :comments org + #+begin_src c :comments org :tangle (eval h_private_func) :noweb yes :exports none qmckl_exit_code qmckl_compute_factor_en ( - const qmckl_context context, - const int64_t walk_num, - const int64_t elec_num, - const int64_t nucl_num, - const int64_t type_nucl_num, - const int64_t* type_nucl_vector, - const int64_t aord_num, - const double* aord_vector, - const double* en_distance_rescaled, - double* const factor_en ); + const qmckl_context context, + const int64_t walk_num, + const int64_t elec_num, + const int64_t nucl_num, + const int64_t type_nucl_num, + const int64_t* type_nucl_vector, + const int64_t aord_num, + const double* aord_vector, + const double* en_distance_rescaled, + double* const factor_en ); #+end_src @@ -2970,22 +2934,21 @@ integer function qmckl_compute_factor_en_deriv_e_f( & end function qmckl_compute_factor_en_deriv_e_f #+end_src - #+CALL: generate_c_header(table=qmckl_factor_en_deriv_e_args,rettyp=get_value("CRetType"),fname=get_value("Name")) +# #+CALL: generate_c_header(table=qmckl_factor_en_deriv_e_args,rettyp=get_value("CRetType"),fname=get_value("Name")) - #+RESULTS: - #+begin_src c :tangle (eval h_func) :comments org + #+begin_src c :comments org :tangle (eval h_private_func) :noweb yes :exports none qmckl_exit_code qmckl_compute_factor_en_deriv_e ( - const qmckl_context context, - const int64_t walk_num, - const int64_t elec_num, - const int64_t nucl_num, - const int64_t type_nucl_num, - const int64_t* type_nucl_vector, - const int64_t aord_num, - const double* aord_vector, - const double* en_distance_rescaled, - const double* en_distance_rescaled_deriv_e, - double* const factor_en_deriv_e ); + const qmckl_context context, + const int64_t walk_num, + const int64_t elec_num, + const int64_t nucl_num, + const int64_t type_nucl_num, + const int64_t* type_nucl_vector, + const int64_t aord_num, + const double* aord_vector, + const double* en_distance_rescaled, + const double* en_distance_rescaled_deriv_e, + double* const factor_en_deriv_e ); #+end_src @@ -2994,7 +2957,7 @@ end function qmckl_compute_factor_en_deriv_e_f #+RESULTS: #+begin_src f90 :tangle (eval f) :comments org :exports none integer(c_int32_t) function qmckl_compute_factor_en_deriv_e & - (context, & + (context, & walk_num, & elec_num, & nucl_num, & @@ -3005,7 +2968,7 @@ end function qmckl_compute_factor_en_deriv_e_f en_distance_rescaled, & en_distance_rescaled_deriv_e, & factor_en_deriv_e) & - bind(C) result(info) + bind(C) result(info) use, intrinsic :: iso_c_binding implicit none @@ -3017,14 +2980,14 @@ end function qmckl_compute_factor_en_deriv_e_f integer (c_int64_t) , intent(in) , value :: type_nucl_num integer (c_int64_t) , intent(in) :: type_nucl_vector(nucl_num) integer (c_int64_t) , intent(in) , value :: aord_num - real (c_double ) , intent(in) :: aord_vector(aord_num + 1, type_nucl_num) + real (c_double ) , intent(in) :: aord_vector(type_nucl_num,aord_num+1) real (c_double ) , intent(in) :: en_distance_rescaled(elec_num,nucl_num,walk_num) real (c_double ) , intent(in) :: en_distance_rescaled_deriv_e(elec_num,nucl_num,4,walk_num) real (c_double ) , intent(out) :: factor_en_deriv_e(elec_num,4,walk_num) integer(c_int32_t), external :: qmckl_compute_factor_en_deriv_e_f info = qmckl_compute_factor_en_deriv_e_f & - (context, & + (context, & walk_num, & elec_num, & nucl_num, & @@ -3366,18 +3329,17 @@ integer function qmckl_compute_een_rescaled_e_f( & end function qmckl_compute_een_rescaled_e_f #+end_src - #+CALL: generate_c_header(table=qmckl_factor_een_rescaled_e_args,rettyp=get_value("CRetType"),fname=get_value("Name")) +# #+CALL: generate_c_header(table=qmckl_factor_een_rescaled_e_args,rettyp=get_value("CRetType"),fname=get_value("Name")) - #+RESULTS: - #+begin_src c :tangle (eval h_func) :comments org + #+begin_src c :comments org :tangle (eval h_private_func) :noweb yes :exports none qmckl_exit_code qmckl_compute_een_rescaled_e ( - const qmckl_context context, - const int64_t walk_num, - const int64_t elec_num, - const int64_t cord_num, - const double rescale_factor_kappa_ee, - const double* ee_distance, - double* const een_rescaled_e ); + const qmckl_context context, + const int64_t walk_num, + const int64_t elec_num, + const int64_t cord_num, + const double rescale_factor_kappa_ee, + const double* ee_distance, + double* const een_rescaled_e ); #+end_src #+CALL: generate_c_interface(table=qmckl_factor_een_rescaled_e_args,rettyp=get_value("CRetType"),fname=get_value("Name")) @@ -3385,9 +3347,8 @@ end function qmckl_compute_een_rescaled_e_f #+RESULTS: #+begin_src f90 :tangle (eval f) :comments org :exports none integer(c_int32_t) function qmckl_compute_een_rescaled_e & - (context, walk_num, elec_num, cord_num, rescale_factor_kappa_ee, & - ee_distance, een_rescaled_e) & - bind(C) result(info) + (context, walk_num, elec_num, cord_num, rescale_factor_kappa_ee, ee_distance, een_rescaled_e) & + bind(C) result(info) use, intrinsic :: iso_c_binding implicit none @@ -3402,7 +3363,7 @@ end function qmckl_compute_een_rescaled_e_f integer(c_int32_t), external :: qmckl_compute_een_rescaled_e_f info = qmckl_compute_een_rescaled_e_f & - (context, walk_num, elec_num, cord_num, rescale_factor_kappa_ee, ee_distance, een_rescaled_e) + (context, walk_num, elec_num, cord_num, rescale_factor_kappa_ee, ee_distance, een_rescaled_e) end function qmckl_compute_een_rescaled_e #+end_src @@ -3597,7 +3558,7 @@ qmckl_exit_code qmckl_provide_een_rescaled_e_deriv_e(qmckl_context context) *** Compute :PROPERTIES: - :Name: qmckl_compute_een_rescaled_e_deriv_e + :Name: qmckl_compute_factor_een_rescaled_e_deriv_e :CRetType: qmckl_exit_code :FRetType: qmckl_exit_code :END: @@ -3704,21 +3665,20 @@ integer function qmckl_compute_factor_een_rescaled_e_deriv_e_f( & end function qmckl_compute_factor_een_rescaled_e_deriv_e_f #+end_src - #+CALL: generate_c_header(table=qmckl_factor_een_rescaled_e_deriv_e_args,rettyp=get_value("CRetType"),fname=get_value("Name")) +# #+CALL: generate_c_header(table=qmckl_factor_een_rescaled_e_deriv_e_args,rettyp=get_value("CRetType"),fname=get_value("Name")) - #+RESULTS: - #+begin_src c :tangle (eval h_func) :comments org - qmckl_exit_code qmckl_compute_factor_een_rescaled_e_deriv_e ( - const qmckl_context context, - const int64_t walk_num, - const int64_t elec_num, - const int64_t cord_num, - const double rescale_factor_kappa_ee, - const double* coord_new, - const double* ee_distance, - const double* een_rescaled_e, - double* const een_rescaled_e_deriv_e ); - #+end_src + #+begin_src c :comments org :tangle (eval h_private_func) :noweb yes :exports none + qmckl_exit_code qmckl_compute_factor_een_rescaled_e_deriv_e ( + const qmckl_context context, + const int64_t walk_num, + const int64_t elec_num, + const int64_t cord_num, + const double rescale_factor_kappa_ee, + const double* coord_new, + const double* ee_distance, + const double* een_rescaled_e, + double* const een_rescaled_e_deriv_e ); + #+end_src #+CALL: generate_c_interface(table=qmckl_factor_een_rescaled_e_deriv_e_args,rettyp=get_value("CRetType"),fname=get_value("Name")) @@ -3726,7 +3686,7 @@ end function qmckl_compute_factor_een_rescaled_e_deriv_e_f #+RESULTS: #+begin_src f90 :tangle (eval f) :comments org :exports none integer(c_int32_t) function qmckl_compute_factor_een_rescaled_e_deriv_e & - (context, & + (context, & walk_num, & elec_num, & cord_num, & @@ -3735,7 +3695,7 @@ end function qmckl_compute_factor_een_rescaled_e_deriv_e_f ee_distance, & een_rescaled_e, & een_rescaled_e_deriv_e) & - bind(C) result(info) + bind(C) result(info) use, intrinsic :: iso_c_binding implicit none @@ -3752,7 +3712,7 @@ end function qmckl_compute_factor_een_rescaled_e_deriv_e_f integer(c_int32_t), external :: qmckl_compute_factor_een_rescaled_e_deriv_e_f info = qmckl_compute_factor_een_rescaled_e_deriv_e_f & - (context, & + (context, & walk_num, & elec_num, & cord_num, & @@ -4126,19 +4086,18 @@ qmckl_exit_code qmckl_compute_een_rescaled_n ( } #+end_src - #+CALL: generate_c_header(table=qmckl_factor_een_rescaled_n_args,rettyp=get_value("CRetType"),fname=get_value("Name")) +# #+CALL: generate_c_header(table=qmckl_factor_een_rescaled_n_args,rettyp=get_value("CRetType"),fname=get_value("Name")) - #+RESULTS: - #+begin_src c :tangle (eval h_func) :comments org + #+begin_src c :comments org :tangle (eval h_private_func) :noweb yes :exports none qmckl_exit_code qmckl_compute_een_rescaled_n ( - const qmckl_context context, - const int64_t walk_num, - const int64_t elec_num, - const int64_t nucl_num, - const int64_t cord_num, - const double rescale_factor_kappa_en, - const double* en_distance, - double* const een_rescaled_n ); + const qmckl_context context, + const int64_t walk_num, + const int64_t elec_num, + const int64_t nucl_num, + const int64_t cord_num, + const double rescale_factor_kappa_en, + const double* en_distance, + double* const een_rescaled_n ); #+end_src *** Test @@ -4433,22 +4392,21 @@ integer function qmckl_compute_factor_een_rescaled_n_deriv_e_f( & end function qmckl_compute_factor_een_rescaled_n_deriv_e_f #+end_src - #+CALL: generate_c_header(table=qmckl_compute_factor_een_rescaled_n_deriv_e_args,rettyp=get_value("CRetType"),fname=get_value("Name")) +# #+CALL: generate_c_header(table=qmckl_compute_factor_een_rescaled_n_deriv_e_args,rettyp=get_value("CRetType"),fname=get_value("Name")) - #+RESULTS: - #+begin_src c :tangle (eval h_func) :comments org + #+begin_src c :comments org :tangle (eval h_private_func) :noweb yes :exports none qmckl_exit_code qmckl_compute_factor_een_rescaled_n_deriv_e ( - const qmckl_context context, - const int64_t walk_num, - const int64_t elec_num, - const int64_t nucl_num, - const int64_t cord_num, - const double rescale_factor_kappa_en, - const double* coord_new, - const double* coord, - const double* en_distance, - const double* een_rescaled_n, - double* const een_rescaled_n_deriv_e ); + const qmckl_context context, + const int64_t walk_num, + const int64_t elec_num, + const int64_t nucl_num, + const int64_t cord_num, + const double rescale_factor_kappa_en, + const double* coord_new, + const double* coord, + const double* en_distance, + const double* een_rescaled_n, + double* const een_rescaled_n_deriv_e ); #+end_src #+CALL: generate_c_interface(table=qmckl_compute_factor_een_rescaled_n_deriv_e_args,rettyp=get_value("CRetType"),fname=get_value("Name")) @@ -4456,7 +4414,7 @@ end function qmckl_compute_factor_een_rescaled_n_deriv_e_f #+RESULTS: #+begin_src f90 :tangle (eval f) :comments org :exports none integer(c_int32_t) function qmckl_compute_factor_een_rescaled_n_deriv_e & - (context, & + (context, & walk_num, & elec_num, & nucl_num, & @@ -4467,7 +4425,7 @@ end function qmckl_compute_factor_een_rescaled_n_deriv_e_f en_distance, & een_rescaled_n, & een_rescaled_n_deriv_e) & - bind(C) result(info) + bind(C) result(info) use, intrinsic :: iso_c_binding implicit none @@ -4481,12 +4439,12 @@ end function qmckl_compute_factor_een_rescaled_n_deriv_e_f real (c_double ) , intent(in) :: coord_new(elec_num,3,walk_num) real (c_double ) , intent(in) :: coord(nucl_num,3) real (c_double ) , intent(in) :: en_distance(nucl_num,elec_num,walk_num) - real (c_double ) , intent(in) :: een_rescaled_n(0:cord_num,nucl_num,elec_num,walk_num) + real (c_double ) , intent(in) :: een_rescaled_n(elec_num,nucl_num,0:cord_num,walk_num) real (c_double ) , intent(out) :: een_rescaled_n_deriv_e(elec_num,4,nucl_num,0:cord_num,walk_num) integer(c_int32_t), external :: qmckl_compute_factor_een_rescaled_n_deriv_e_f info = qmckl_compute_factor_een_rescaled_n_deriv_e_f & - (context, & + (context, & walk_num, & elec_num, & nucl_num, & @@ -4891,64 +4849,54 @@ qmckl_exit_code qmckl_provide_tmp_c(qmckl_context context) } /* Choose the correct compute function (depending on offload type) */ - switch(ctx->jastrow.offload_type) { - case OFFLOAD_OPENACC: - #ifdef HAVE_OPENACC_OFFLOAD - rc = - qmckl_compute_tmp_c_acc_offload(context, - ctx->jastrow.cord_num, - ctx->electron.num, - ctx->nucleus.num, - ctx->electron.walk_num, - ctx->jastrow.een_rescaled_e, - ctx->jastrow.een_rescaled_n, - ctx->jastrow.tmp_c); - #else - rc = qmckl_compute_tmp_c(context, - ctx->jastrow.cord_num, - ctx->electron.num, - ctx->nucleus.num, - ctx->electron.walk_num, - ctx->jastrow.een_rescaled_e, - ctx->jastrow.een_rescaled_n, - ctx->jastrow.tmp_c); - - #endif - break; - case OFFLOAD_CUBLAS: - #ifdef HAVE_CUBLAS_OFFLOAD - rc = - qmckl_compute_tmp_c_cublas_offload(context, - ctx->jastrow.cord_num, - ctx->electron.num, - ctx->nucleus.num, - ctx->electron.walk_num, - ctx->jastrow.een_rescaled_e, - ctx->jastrow.een_rescaled_n, - ctx->jastrow.tmp_c); - #else - rc = qmckl_compute_tmp_c(context, - ctx->jastrow.cord_num, - ctx->electron.num, - ctx->nucleus.num, - ctx->electron.walk_num, - ctx->jastrow.een_rescaled_e, - ctx->jastrow.een_rescaled_n, - ctx->jastrow.tmp_c); - #endif - break; - default: - rc = qmckl_compute_tmp_c(context, - ctx->jastrow.cord_num, - ctx->electron.num, - ctx->nucleus.num, - ctx->electron.walk_num, - ctx->jastrow.een_rescaled_e, - ctx->jastrow.een_rescaled_n, - ctx->jastrow.tmp_c); - break; +#ifdef HAVE_HPC + const bool gpu_offload = ctx->jastrow.gpu_offload; +#else + const bool gpu_offload = false; +#endif + + if (gpu_offload) { +#ifdef HAVE_CUBLAS_OFFLOAD + rc = qmckl_compute_tmp_c_cublas_offload(context, + ctx->jastrow.cord_num, + ctx->electron.num, + ctx->nucleus.num, + ctx->electron.walk_num, + ctx->jastrow.een_rescaled_e, + ctx->jastrow.een_rescaled_n, + ctx->jastrow.tmp_c); +#elif HAVE_OPENACC_OFFLOAD + rc = qmckl_compute_tmp_c_acc_offload(context, + ctx->jastrow.cord_num, + ctx->electron.num, + ctx->nucleus.num, + ctx->electron.walk_num, + ctx->jastrow.een_rescaled_e, + ctx->jastrow.een_rescaled_n, + ctx->jastrow.tmp_c); +#elif HAVE_OPENMP_OFFLOAD + rc = qmckl_compute_tmp_c_omp_offload(context, + ctx->jastrow.cord_num, + ctx->electron.num, + ctx->nucleus.num, + ctx->electron.walk_num, + ctx->jastrow.een_rescaled_e, + ctx->jastrow.een_rescaled_n, + ctx->jastrow.tmp_c); +#else + rc = QMCKL_FAILURE; +#endif + } else { + rc = qmckl_compute_tmp_c(context, + ctx->jastrow.cord_num, + ctx->electron.num, + ctx->nucleus.num, + ctx->electron.walk_num, + ctx->jastrow.een_rescaled_e, + ctx->jastrow.een_rescaled_n, + ctx->jastrow.tmp_c); } - + ctx->jastrow.tmp_c_date = ctx->date; } @@ -4988,18 +4936,44 @@ qmckl_exit_code qmckl_provide_dtmp_c(qmckl_context context) ctx->jastrow.dtmp_c = dtmp_c; } - switch(ctx->jastrow.offload_type) { - case OFFLOAD_OPENACC: - #ifdef HAVE_OPENACC_OFFLOAD - rc = qmckl_compute_dtmp_c_acc_offload(context, - ctx->jastrow.cord_num, - ctx->electron.num, - ctx->nucleus.num, - ctx->electron.walk_num, - ctx->jastrow.een_rescaled_e_deriv_e, - ctx->jastrow.een_rescaled_n, - ctx->jastrow.dtmp_c); - #else +#ifdef HAVE_HPC + const bool gpu_offload = ctx->jastrow.gpu_offload; +#else + const bool gpu_offload = false; +#endif + + if (gpu_offload) { +#ifdef HAVE_CUBLAS_OFFLOAD + rc = qmckl_compute_dtmp_c_cublas_offload(context, + ctx->jastrow.cord_num, + ctx->electron.num, + ctx->nucleus.num, + ctx->electron.walk_num, + ctx->jastrow.een_rescaled_e_deriv_e, + ctx->jastrow.een_rescaled_n, + ctx->jastrow.dtmp_c); +#elif HAVE_OPENACC_OFFLOAD + rc = qmckl_compute_dtmp_c_acc_offload(context, + ctx->jastrow.cord_num, + ctx->electron.num, + ctx->nucleus.num, + ctx->electron.walk_num, + ctx->jastrow.een_rescaled_e_deriv_e, + ctx->jastrow.een_rescaled_n, + ctx->jastrow.dtmp_c); +#elif HAVE_OPENMP_OFFLOAD + rc = qmckl_compute_dtmp_c_omp_offload(context, + ctx->jastrow.cord_num, + ctx->electron.num, + ctx->nucleus.num, + ctx->electron.walk_num, + ctx->jastrow.een_rescaled_e_deriv_e, + ctx->jastrow.een_rescaled_n, + ctx->jastrow.dtmp_c); +#else + rc = QMCKL_FAILURE; +#endif + } else { rc = qmckl_compute_dtmp_c(context, ctx->jastrow.cord_num, ctx->electron.num, @@ -5008,39 +4982,6 @@ qmckl_exit_code qmckl_provide_dtmp_c(qmckl_context context) ctx->jastrow.een_rescaled_e_deriv_e, ctx->jastrow.een_rescaled_n, ctx->jastrow.dtmp_c); - #endif - break; - case OFFLOAD_CUBLAS: - #ifdef HAVE_CUBLAS_OFFLOAD - rc = qmckl_compute_dtmp_c_acc_offload(context, - ctx->jastrow.cord_num, - ctx->electron.num, - ctx->nucleus.num, - ctx->electron.walk_num, - ctx->jastrow.een_rescaled_e_deriv_e, - ctx->jastrow.een_rescaled_n, - ctx->jastrow.dtmp_c); - #else - rc = qmckl_compute_dtmp_c(context, - ctx->jastrow.cord_num, - ctx->electron.num, - ctx->nucleus.num, - ctx->electron.walk_num, - ctx->jastrow.een_rescaled_e_deriv_e, - ctx->jastrow.een_rescaled_n, - ctx->jastrow.dtmp_c); - #endif - break; - default: - rc = qmckl_compute_dtmp_c(context, - ctx->jastrow.cord_num, - ctx->electron.num, - ctx->nucleus.num, - ctx->electron.walk_num, - ctx->jastrow.een_rescaled_e_deriv_e, - ctx->jastrow.een_rescaled_n, - ctx->jastrow.dtmp_c); - break; } if (rc != QMCKL_SUCCESS) { @@ -5148,14 +5089,13 @@ qmckl_exit_code qmckl_compute_dim_cord_vect ( } #+end_src - #+CALL: generate_c_header(table=qmckl_factor_dim_cord_vect_args,rettyp=get_value("CRetType"),fname=get_value("Name")) +# #+CALL: generate_c_header(table=qmckl_factor_dim_cord_vect_args,rettyp=get_value("CRetType"),fname=get_value("Name")) - #+RESULTS: - #+begin_src c :tangle (eval h_func) :comments org + #+begin_src c :comments org :tangle (eval h_private_func) :noweb yes :exports none qmckl_exit_code qmckl_compute_dim_cord_vect ( - const qmckl_context context, - const int64_t cord_num, - int64_t* const dim_cord_vect ); + const qmckl_context context, + const int64_t cord_num, + int64_t* const dim_cord_vect ); #+end_src @@ -5224,18 +5164,17 @@ integer function qmckl_compute_cord_vect_full_f( & end function qmckl_compute_cord_vect_full_f #+end_src - #+CALL: generate_c_header(table=qmckl_factor_cord_vect_full_args,rettyp=get_value("CRetType"),fname=get_value("Name")) +# #+CALL: generate_c_header(table=qmckl_factor_cord_vect_full_args,rettyp=get_value("CRetType"),fname=get_value("Name")) - #+RESULTS: - #+begin_src c :tangle (eval h_func) :comments org + #+begin_src c :comments org :tangle (eval h_private_func) :noweb yes :exports none qmckl_exit_code qmckl_compute_cord_vect_full ( - const qmckl_context context, - const int64_t nucl_num, - const int64_t dim_cord_vect, - const int64_t type_nucl_num, - const int64_t* type_nucl_vector, - const double* cord_vector, - double* const cord_vect_full ); + const qmckl_context context, + const int64_t nucl_num, + const int64_t dim_cord_vect, + const int64_t type_nucl_num, + const int64_t* type_nucl_vector, + const double* cord_vector, + double* const cord_vect_full ); #+end_src @@ -5244,8 +5183,8 @@ end function qmckl_compute_cord_vect_full_f #+RESULTS: #+begin_src f90 :tangle (eval f) :comments org :exports none integer(c_int32_t) function qmckl_compute_cord_vect_full & - (context, nucl_num, dim_cord_vect, type_nucl_num, type_nucl_vector, cord_vector, cord_vect_full) & - bind(C) result(info) + (context, nucl_num, dim_cord_vect, type_nucl_num, type_nucl_vector, cord_vector, cord_vect_full) & + bind(C) result(info) use, intrinsic :: iso_c_binding implicit none @@ -5260,7 +5199,7 @@ end function qmckl_compute_cord_vect_full_f integer(c_int32_t), external :: qmckl_compute_cord_vect_full_f info = qmckl_compute_cord_vect_full_f & - (context, nucl_num, dim_cord_vect, type_nucl_num, type_nucl_vector, cord_vector, cord_vect_full) + (context, nucl_num, dim_cord_vect, type_nucl_num, type_nucl_vector, cord_vector, cord_vect_full) end function qmckl_compute_cord_vect_full #+end_src @@ -5381,15 +5320,14 @@ qmckl_exit_code qmckl_compute_lkpm_combined_index ( } #+end_src - #+CALL: generate_c_header(table=qmckl_factor_lkpm_combined_index_args,rettyp=get_value("CRetType"),fname=get_value("Name")) +# #+CALL: generate_c_header(table=qmckl_factor_lkpm_combined_index_args,rettyp=get_value("CRetType"),fname=get_value("Name")) - #+RESULTS: - #+begin_src c :tangle (eval h_func) :comments org + #+begin_src c :comments org :tangle (eval h_private_func) :noweb yes :exports none qmckl_exit_code qmckl_compute_lkpm_combined_index ( - const qmckl_context context, - const int64_t cord_num, - const int64_t dim_cord_vect, - int64_t* const lkpm_combined_index ); + const qmckl_context context, + const int64_t cord_num, + const int64_t dim_cord_vect, + int64_t* const lkpm_combined_index ); #+end_src @@ -5413,6 +5351,38 @@ qmckl_exit_code qmckl_compute_lkpm_combined_index ( | ~een_rescaled_n~ | ~double[walk_num][0:cord_num][nucl_num][elec_num]~ | in | Electron-nucleus rescaled factor | | ~tmp_c~ | ~double[walk_num][0:cord_num-1][0:cord_num][nucl_num][elec_num]~ | out | vector of non-zero coefficients | + #+begin_src c :comments org :tangle (eval c) :noweb yes +qmckl_exit_code qmckl_compute_tmp_c (const qmckl_context context, + const int64_t cord_num, + const int64_t elec_num, + const int64_t nucl_num, + const int64_t walk_num, + const double* een_rescaled_e, + const double* een_rescaled_n, + double* const tmp_c ) +{ +#ifdef HAVE_HPC + return qmckl_compute_tmp_c_hpc(context, cord_num, elec_num, nucl_num, walk_num, een_rescaled_e, een_rescaled_n, tmp_c); +#else + return qmckl_compute_tmp_c_doc(context, cord_num, elec_num, nucl_num, walk_num, een_rescaled_e, een_rescaled_n, tmp_c); +#endif +} + #+end_src + +# #+CALL: generate_c_header(table=qmckl_factor_tmp_c_args,rettyp=get_value("CRetType"),fname="qmckl_compute_tmp_c") + + #+begin_src c :comments org :tangle (eval h_private_func) :noweb yes :exports none + qmckl_exit_code qmckl_compute_tmp_c ( + const qmckl_context context, + const int64_t cord_num, + const int64_t elec_num, + const int64_t nucl_num, + const int64_t walk_num, + const double* een_rescaled_e, + const double* een_rescaled_n, + double* const tmp_c ); + #+end_src + #+begin_src f90 :comments org :tangle (eval f) :noweb yes integer function qmckl_compute_tmp_c_doc_f( & context, cord_num, elec_num, nucl_num, & @@ -5481,8 +5451,20 @@ integer function qmckl_compute_tmp_c_doc_f( & end function qmckl_compute_tmp_c_doc_f #+end_src -#+CALL: generate_c_interface(table=qmckl_factor_tmp_c_args,rettyp=get_value("FRetType"),fname="qmckl_compute_tmp_c_doc") + #+begin_src c :tangle (eval h_private_func) :comments org +qmckl_exit_code qmckl_compute_tmp_c_doc ( + const qmckl_context context, + const int64_t cord_num, + const int64_t elec_num, + const int64_t nucl_num, + const int64_t walk_num, + const double* een_rescaled_e, + const double* een_rescaled_n, + double* const tmp_c ); + #+end_src + #+CALL: generate_c_interface(table=qmckl_factor_tmp_c_args,rettyp=get_value("FRetType"),fname="qmckl_compute_tmp_c_doc") + #+RESULTS: #+begin_src f90 :tangle (eval f) :comments org :exports none integer(c_int32_t) function qmckl_compute_tmp_c_doc & @@ -5508,6 +5490,7 @@ integer(c_int32_t) function qmckl_compute_tmp_c_doc & end function qmckl_compute_tmp_c_doc #+end_src +**** CPU :noexport: #+begin_src c :comments org :tangle (eval c) :noweb yes qmckl_exit_code qmckl_compute_tmp_c_hpc ( @@ -5559,16 +5542,15 @@ qmckl_exit_code qmckl_compute_tmp_c_hpc ( const int64_t bf = elec_num*nucl_num*(cord_num+1); const int64_t cf = bf; +#ifdef HAVE_OPENMP +#pragma omp parallel for collapse(2) +#endif for (int64_t nw=0; nw < walk_num; ++nw) { for (int64_t i=0; i