More fixups.

Author: olupton

.gitlab-ci.yml

@@ -93,7 +93,7 @@ build:coreneuron+nmodl:gpu:
     SPACK_PACKAGE: coreneuron
     # +report pulls in a lot of dependencies and the tests fail.
     # See https://github.com/BlueBrain/CoreNeuron/issues/518 re: build_type
-    SPACK_PACKAGE_SPEC: +nmodl+openmp+gpu+tests~legacy-unit~report build_type=RelWithDebInfo
+    SPACK_PACKAGE_SPEC: +nmodl+openmp+gpu+tests~legacy-unit~report+sympy build_type=RelWithDebInfo
   extends:
     - .spack_build
     - .spack_nvhpc
@@ -104,7 +104,8 @@ build:coreneuron+nmodl~openmp:gpu:
     SPACK_PACKAGE: coreneuron
     # +report pulls in a lot of dependencies and the tests fail.
     # See https://github.com/BlueBrain/CoreNeuron/issues/518 re: build_type
-    SPACK_PACKAGE_SPEC: +nmodl~openmp+gpu+tests~legacy-unit~report build_type=RelWithDebInfo
+    # Sympy + OpenMP target offload does not currently work with NVHPC
+    SPACK_PACKAGE_SPEC: +nmodl~openmp+gpu+tests~legacy-unit~report~sympy build_type=RelWithDebInfo
   extends:
     - .spack_build
     - .spack_nvhpc
@@ -178,7 +179,7 @@ build:neuron+nmodl~openmp:gpu:
     # Build py-cython and py-numpy with GCC instead of NVHPC.
     - SPACK_PACKAGE_DEPENDENCIES="${SPACK_PACKAGE_DEPENDENCIES}^py-cython%gcc^py-numpy%gcc"
     - !reference [.spack_build, before_script]
-  needs: ["build:coreneuron+nmodl:gpu"]
+  needs: ["build:coreneuron+nmodl~openmp:gpu"]
 
 build:neuron:gpu:
   stage: build_neuron

coreneuron/mechanism/mech/mod2c_core_thread.hpp

@@ -100,7 +100,9 @@ extern void* nrn_cons_sparseobj(SPFUN, int, Memb_list*, _threadargsproto_);
 extern void _nrn_destroy_sparseobj_thread(SparseObj* so);
 
 nrn_pragma_acc(routine seq)
+nrn_pragma_omp(declare target)
 extern int nrn_kinetic_steer(int, SparseObj*, double*, _threadargsproto_);
+nrn_pragma_omp(end declare target)
 #define spfun(arg1, arg2, arg3) nrn_kinetic_steer(arg1, arg2, arg3, _threadargs_);
 
 // derived from nrn/src/scopmath/euler.c
@@ -119,6 +121,7 @@ static inline int euler_thread(int neqn, int* var, int* der, DIFUN fun, _threada
     return 0;
 }
 
+nrn_pragma_omp(declare target)
 nrn_pragma_acc(routine seq)
 extern int derivimplicit_thread(int, int*, int*, DIFUN, _threadargsproto_);
 nrn_pragma_acc(routine seq)
@@ -141,6 +144,7 @@ nrn_pragma_acc(routine seq)
 extern double _modl_get_dt_thread(NrnThread*);
 nrn_pragma_acc(routine seq)
 extern void _modl_set_dt_thread(double, NrnThread*);
+nrn_pragma_omp(end declare target)
 
 void nrn_sparseobj_copyto_device(SparseObj* so);
 void nrn_sparseobj_delete_from_device(SparseObj* so);

coreneuron/mechanism/membfunc.hpp

@@ -179,7 +179,9 @@ extern void nrn2ncs_outputevent(int netcon_output_index, double firetime);
 extern bool nrn_use_localgid_;
 extern void net_sem_from_gpu(int sendtype, int i_vdata, int, int ith, int ipnt, double, double);
 nrn_pragma_acc(routine seq)
+nrn_pragma_omp(declare target)
 extern int at_time(NrnThread*, double);
+nrn_pragma_omp(end declare target)
 
 // _OPENACC and/or NET_RECEIVE_BUFFERING
 extern void net_sem_from_gpu(int, int, int, int, int, double, double);

coreneuron/utils/ivocvect.hpp

@@ -9,6 +9,8 @@
 #ifndef ivoc_vector_h
 #define ivoc_vector_h
 
+#include "coreneuron/utils/offload.hpp"
+
 #include <cstdio>
 #include <utility>
 
@@ -76,9 +78,9 @@ extern double* vector_vec(IvocVect* v);
 
 // retro-compatibility API
 extern void* vector_new1(int n);
-#pragma acc routine seq
+nrn_pragma_acc(routine seq)
 extern int vector_capacity(void* v);
-#pragma acc routine seq
+nrn_pragma_acc(routine seq)
 extern double* vector_vec(void* v);
 
 }  // namespace coreneuron

coreneuron/utils/randoms/nrnran123.h

@@ -37,6 +37,8 @@ of the full distribution available from
 #define R123_USE_GNU_UINT128            1
 #endif
 
+#include "coreneuron/utils/offload.hpp"
+
 #include <Random123/philox.h>
 #include <inttypes.h>
 
@@ -46,12 +48,7 @@ of the full distribution available from
 #define CORENRN_HOST_DEVICE
 #endif
 
-// Is there actually any harm leaving the pragma in when DISABLE_OPENACC is true?
-#if defined(_OPENACC) && !defined(DISABLE_OPENACC)
-#define CORENRN_HOST_DEVICE_ACC CORENRN_HOST_DEVICE _Pragma("acc routine seq")
-#else
-#define CORENRN_HOST_DEVICE_ACC CORENRN_HOST_DEVICE
-#endif
+#define CORENRN_HOST_DEVICE_ACC CORENRN_HOST_DEVICE nrn_pragma_acc(routine seq)
 
 // Some files are compiled with DISABLE_OPENACC, and some builds have no GPU
 // support at all. In these two cases, request that the random123 state is
@@ -100,6 +97,7 @@ void nrnran123_deletestream(nrnran123_State* s,
                             bool use_unified_memory = CORENRN_RAN123_USE_UNIFIED_MEMORY);
 
 /* minimal data stream */
+nrn_pragma_omp(declare target)
 CORENRN_HOST_DEVICE_ACC void nrnran123_getseq(nrnran123_State*, uint32_t* seq, char* which);
 CORENRN_HOST_DEVICE_ACC void nrnran123_getids(nrnran123_State*, uint32_t* id1, uint32_t* id2);
 CORENRN_HOST_DEVICE_ACC void nrnran123_getids3(nrnran123_State*,
@@ -128,6 +126,7 @@ CORENRN_HOST_DEVICE_ACC nrnran123_array4x32 nrnran123_iran(uint32_t seq,
                                                            uint32_t id1,
                                                            uint32_t id2);
 CORENRN_HOST_DEVICE_ACC double nrnran123_uint2dbl(uint32_t);
+nrn_pragma_omp(end declare target)
 }  // namespace coreneuron
 
 #endif