Fix GPU compilation issue when nrn_ghk is used in the MOD files (#498)

 - nrn_ghk was recently added but `acc routine seq` was missing
 - nrn_ghk uses celsius which is an extern variable. All extern
   variables require `acc declare create/copyin` caluse.
 - As global variables can be initialized after device attachment
   (celsius in this case), we need to make sure celsius is updated
   on the device.
 - Replacing `data` clause with `update device` is sufficient. In
   this case `data` clause doesn't update value on device when
   checked in eion.cpp.

Author: pramodk

coreneuron/apps/main1.cpp

@@ -490,7 +490,7 @@ extern "C" int run_solve_core(int argc, char** argv) {
 
     // clang-format off
 
-    #pragma acc data copyin(celsius, secondorder) if (compute_gpu)
+    #pragma acc update device(celsius, secondorder) if (compute_gpu)
     // clang-format on
     {
         double v = corenrn_param.voltage;

coreneuron/mechanism/eion.cpp

@@ -189,6 +189,10 @@ the USEION statement of any model using this ion\n",
 #define FARADAY     _faraday_codata2018
 #define gasconstant _gasconstant_codata2018
 #endif
+
+// extern variables require acc declare
+#pragma acc declare create(celsius)
+
 #define ktf (1000. * gasconstant * (celsius + 273.15) / FARADAY)
 
 double nrn_nernst(double ci, double co, double z, double celsius) {

coreneuron/mechanism/membfunc.hpp

@@ -110,6 +110,8 @@ extern void nrn_writes_conc(int, int);
 extern void nrn_wrote_conc(int, double*, int, int, double**, double, int);
 #pragma acc routine seq
 double nrn_nernst(double ci, double co, double z, double celsius);
+#pragma acc routine seq
+extern double nrn_ghk(double v, double ci, double co, double z);
 extern void hoc_register_prop_size(int, int, int);
 extern void hoc_register_dparam_semantics(int type, int, const char* name);