OpenACC pragma annotation was missing for kernels using euler (#85)

* cudaMemset wrong pointer fix
* User provided flags should take precedence

Author: pramodk

CMakeLists.txt

@@ -206,8 +206,8 @@ if(ENABLE_OPENACC)
         if (ENABLE_OPENACC_INFO)
             set(ACC_FLAGS "${ACC_FLAGS} -Minfo=acc")
         endif()
-        set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${ACC_FLAGS}")
-        set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${ACC_FLAGS}")
+        set(CMAKE_C_FLAGS "${ACC_FLAGS} ${CMAKE_C_FLAGS}")
+        set(CMAKE_CXX_FLAGS "${ACC_FLAGS} ${CMAKE_CXX_FLAGS}")
     else()
         message(WARNING "OpenACC implementation is only supported and tested using only PGI")
         message(WARNING "Add required compiler flags to enabled OpenACC")

coreneuron/kinderiv.py

@@ -57,6 +57,10 @@
   fout.write('#pragma acc routine seq\n')
   fout.write('extern int %s%s(void*, double*, _threadargsproto_);\n' % (item[0], item[1]))
 
+for item in euler:
+  fout.write('#pragma acc routine seq\n')
+  fout.write('extern int %s%s(_threadargsproto_);\n' % (item[0], item[1]))
+
 fout.write("\n/* callback indices */\n")
 derivoffset = 1
 kinoffset = 1

coreneuron/utils/randoms/nrnran123.cu

@@ -149,7 +149,7 @@ nrnran123_State* nrnran123_newstream3(uint32_t id1, uint32_t id2, uint32_t id3)
     nrnran123_State* s;
 
     cudaMalloc((void**)&s, sizeof(nrnran123_State));
-    cudaMemset((void**)&s, 0, sizeof(nrnran123_State));
+    cudaMemset((void*)s, 0, sizeof(nrnran123_State));
 
     nrnran123_setup_cuda_newstream<<<1, 1>>>(s, id1, id2, id3);
     cudaDeviceSynchronize();