Bug fix in gpu implementation : global_vars.cpp was compiled with (#83)

OpenACC resulting in nrnran123_set_globalindex call for gpu device
kernel. This device call was resulting in failure and hence abnormal
program termination.

* update mod2c submodule

Author: pramodk

coreneuron/CMakeLists.txt

@@ -235,7 +235,13 @@ endif()
 set(GENERATED_MECH_C_FILES ${MOD_FUNC_C} ${MOD_FUNC_PTRS_C} ${MOD2C_STDMECH_OUTPUTS} ${MOD2C_OPTMECH_OUTPUTS})
 
 # artificial cells must be on cpu, defaul nrnran123.c is for cpu, nrn_setup.cpp uses nrnran123 for only memory calculation purpose which should use cpu version of nrnran123
-set(NOACC_MECH_C_FILES ${CMAKE_CURRENT_BINARY_DIR}/netstim.c ${CMAKE_CURRENT_BINARY_DIR}/netstim_inhpoisson.c ${CMAKE_CURRENT_BINARY_DIR}/pattern.c ${CMAKE_CURRENT_SOURCE_DIR}/utils/randoms/nrnran123.c ${CMAKE_CURRENT_SOURCE_DIR}/nrniv/nrn_setup.cpp)
+set(NOACC_MECH_C_FILES
+    ${CMAKE_CURRENT_BINARY_DIR}/netstim.c
+    ${CMAKE_CURRENT_BINARY_DIR}/netstim_inhpoisson.c
+    ${CMAKE_CURRENT_BINARY_DIR}/pattern.c
+    ${CMAKE_CURRENT_SOURCE_DIR}/utils/randoms/nrnran123.c
+    ${CMAKE_CURRENT_SOURCE_DIR}/nrniv/nrn_setup.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/nrniv/global_vars.cpp)
 
 if(ENABLE_OPENACC)
     set_source_files_properties(${GENERATED_MECH_C_FILES} PROPERTIES COMPILE_FLAGS "")

coreneuron/utils/randoms/nrnran123.c

@@ -42,21 +42,15 @@ size_t nrnran123_instance_count() {
     return instance_count_;
 }
 
-/* now this is declated in nrnran123.h so that its available for prototype declaration
-
-struct nrnran123_State {
-        philox4x32_ctr_t c;
-        philox4x32_ctr_t r;
-        unsigned char which_;
-};
-*/
-
 size_t nrnran123_state_size() {
     return sizeof(nrnran123_State);
 }
 
 void nrnran123_set_globalindex(uint32_t gix) {
     k.v[0] = gix;
+    #if (defined(__CUDACC__) || defined(_OPENACC))
+    nrnran123_set_gpu_globalindex(gix);
+    #endif
 }
 
 /* if one sets the global, one should reset all the stream sequences. */

coreneuron/utils/randoms/nrnran123.cu

@@ -43,7 +43,7 @@ __device__ size_t nrnran123_state_size() {
     return sizeof(nrnran123_State);
 }
 
-__device__ void nrnran123_set_globalindex(uint32_t gix) {
+__global__ void nrnran123_set_globalindex(uint32_t gix) {
     k.v[0] = gix;
 }
 
@@ -164,3 +164,9 @@ void nrnran123_deletestream(nrnran123_State* s) {
 
     cudaFree(s);
 }
+
+/* set global index for random123 stream on gpu */
+void nrnran123_set_gpu_globalindex(uint32_t gix) {
+    nrnran123_set_globalindex<<<1,1>>>(gix);
+    cudaDeviceSynchronize();
+}

coreneuron/utils/randoms/nrnran123.h

@@ -98,7 +98,7 @@ extern DEVICE void nrnran123_mutconstruct(void);
 
 /* global index. eg. run number */
 /* all generator instances share this global index */
-extern DEVICE void nrnran123_set_globalindex(uint32_t gix);
+extern GLOBAL void nrnran123_set_globalindex(uint32_t gix);
 extern DEVICE uint32_t nrnran123_get_globalindex();
 
 extern DEVICE size_t nrnran123_instance_count(void);
@@ -155,7 +155,7 @@ extern DEVICE double nrnran123_gauss(nrnran123_State*); /* mean 0.0, std 1.0 */
 /* more fundamental (stateless) (though the global index is still used) */
 extern DEVICE nrnran123_array4x32 nrnran123_iran(uint32_t seq, uint32_t id1, uint32_t id2);
 extern DEVICE double nrnran123_uint2dbl(uint32_t);
-
+extern void nrnran123_set_gpu_globalindex(uint32_t gix);
 #if defined(__cplusplus)
 }
 #endif