OSCR

Realistic coupling enables flexible macroscopic traveling waves in the mouse cortex.

Code ↔ Paper

1 match between paragraphs of the paper and lines of its authors' code, computed by the harvester (lexical-v1). Click a colored paragraph or line to see its counterpart.

The 1 match
  1. [1] § Materials and methods › Simulation ↔ code_simulation.zip/coupling_speed_test/ephys.cu, lines 27–123 · score 0.57 · gate variables, 0.02–0.05, simulation, voltage, uniform

Paper

Loaded from Europe PMC by your browser, not stored by OSCR: doi.org · Europe PMC

The paper is loaded when this pane is shown.

The authors' code

CUDA · 403 lines · 19 KB · CC-BY-4.0 · 1 match

  1. #include "ephys.cuh"
  2. # define M_PI 3.14159265358979323846 /* pi */
  3. Ephys::~Ephys()
  4. {
  5. for (Float*& dev_ptr : sim_vars_cpu)
  6. cudaFree(dev_ptr);
  7. }
  9. HHEphys::HHEphys(int const N_)
  10. : Ephys(N_, HH::leapfrog_dt)
  11. {
  12. using namespace HH;
  13. // Set up simulation variables. Stored in a struct of arrays.
  14. sim_vars_cpu.resize(Params::PARAM_COUNT);
  15. thrust::host_vector<Float> cpu_arr(N, 0);
  17. // Zero arrays.
  18. for (Float*& arr : sim_vars_cpu)
  19. {
  20. cudaMalloc(&arr, N * sizeof(Float));
  21. cudaMemcpy(arr, cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  22. }
  23. sim_vars = sim_vars_cpu;
  24. SetVariablesRandomly();
  25. }
  27. void HHEphys::SetVariablesRandomly()
  28. {
  29. using namespace HH;
  30. thrust::host_vector<Float> cpu_arr(N, 0.05);
  32. thrust::fill_n(cpu_arr.begin(), N, g_ampa_white);
  33. cudaMemcpy(sim_vars_cpu[Params::G_AMPA_WHITE], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  35. thrust::fill_n(cpu_arr.begin(), N, g_ampa_grey);
  36. cudaMemcpy(sim_vars_cpu[Params::G_AMPA_GREY], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  38. thrust::fill_n(cpu_arr.begin(), N, g_gaba);
  39. cudaMemcpy(sim_vars_cpu[Params::G_GABA], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  41. // Set voltages randomly
  42. auto& gen{ Utility::GetRandomEngine() };
  43. std::normal_distribution<> voltage_dist(-65,-5);
  44. for (auto& val : cpu_arr)
  45. val = voltage_dist(gen);
  46. cudaMemcpy(sim_vars_cpu[Params::V], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  48. // Set gating variables
  49. std::uniform_real_distribution<Float> gating_var_dist(0.02, 0.05);
  51. for (auto& val : cpu_arr)
  52. val = gating_var_dist(gen);
  53. cudaMemcpy(sim_vars_cpu[Params::M], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  55. for (auto& val : cpu_arr)
  56. val = gating_var_dist(gen);
  57. cudaMemcpy(sim_vars_cpu[Params::H], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  59. for (auto& val : cpu_arr)
  60. val = gating_var_dist(gen);
  61. cudaMemcpy(sim_vars_cpu[Params::N], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  63. for (auto& val : cpu_arr)
  64. val = gating_var_dist(gen);
  65. cudaMemcpy(sim_vars_cpu[Params::P], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  67. for (auto& val : cpu_arr)
  68. val = gating_var_dist(gen);
  69. cudaMemcpy(sim_vars_cpu[Params::Q], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  71. for (auto& val : cpu_arr)
  72. val = gating_var_dist(gen);
  73. cudaMemcpy(sim_vars_cpu[Params::R], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  75. for (auto& val : cpu_arr)
  76. val = gating_var_dist(gen);
  77. cudaMemcpy(sim_vars_cpu[Params::S], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  80. for (auto& val : cpu_arr)
  81. val = gating_var_dist(gen);
  82. cudaMemcpy(sim_vars_cpu[Params::U], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  84. thrust::host_vector<Float> ei_info(N, 0.0);
  85. cudaMemcpy(sim_vars_cpu[Params::EI_INFO], ei_info.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  87. thrust::host_vector<Float> type_info(N, 0.0);
  88. cudaMemcpy(sim_vars_cpu[Params::TYPE_INFO], type_info.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  90. thrust::fill_n(cpu_arr.begin(), N, 0.05);
  91. //cudaMemcpy(sim_vars_cpu[Params::LOCAL_IY], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  92. //cudaMemcpy(sim_vars_cpu[Params::LOCAL_EY], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  93. cudaMemcpy(sim_vars_cpu[Params::IY], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  94. cudaMemcpy(sim_vars_cpu[Params::EY], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  96. // No neurotransmitter input/output or connections
  97. thrust::fill_n(cpu_arr.begin(), N, 0.001);
  98. cudaMemcpy(sim_vars_cpu[Params::LOCAL_IOUTPUT], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  99. cudaMemcpy(sim_vars_cpu[Params::LOCAL_IIN], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  100. cudaMemcpy(sim_vars_cpu[Params::LOCAL_ICOUNT], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  101. cudaMemcpy(sim_vars_cpu[Params::LOCAL_EOUTPUT], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  102. cudaMemcpy(sim_vars_cpu[Params::LOCAL_EIN], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  103. cudaMemcpy(sim_vars_cpu[Params::LOCAL_ECOUNT], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  104. cudaMemcpy(sim_vars_cpu[Params::IOUTPUT], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  105. cudaMemcpy(sim_vars_cpu[Params::IIN], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  106. cudaMemcpy(sim_vars_cpu[Params::ICOUNT], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  107. cudaMemcpy(sim_vars_cpu[Params::EOUTPUT], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  108. cudaMemcpy(sim_vars_cpu[Params::EIN], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  109. cudaMemcpy(sim_vars_cpu[Params::ECOUNT], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  111. // Set applied current
  112. //thrust::fill_n(cpu_arr.begin(), N, appcur);
  113. std::normal_distribution<Float> appcur_dist(appcur, 0.0);
  114. for (auto& val : cpu_arr) {
  115. val = appcur_dist(gen);
  116. }
  118. cudaMemcpy(sim_vars_cpu[Params::APPCUR], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  119. gpuErrchk(cudaDeviceSynchronize());
  121. // Copy pointers from CPU to GPU
  122. sim_vars = sim_vars_cpu;
  123. }
  125. void HHEphys::RandomizeInitialConditionsKeepEI()
  126. {
  127. using namespace HH;
  128. thrust::host_vector<Float> cpu_arr(N, 0.05);
  130. thrust::fill_n(cpu_arr.begin(), N, g_ampa_white);
  131. cudaMemcpy(sim_vars_cpu[Params::G_AMPA_WHITE], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  133. thrust::fill_n(cpu_arr.begin(), N, g_ampa_grey);
  134. cudaMemcpy(sim_vars_cpu[Params::G_AMPA_GREY], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  136. thrust::fill_n(cpu_arr.begin(), N, g_gaba);
  137. cudaMemcpy(sim_vars_cpu[Params::G_GABA], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  139. auto& gen{ Utility::GetRandomEngine() };
  140. std::normal_distribution<> voltage_dist(-65, -5);
  141. for (auto& val : cpu_arr)
  142. val = voltage_dist(gen);
  143. cudaMemcpy(sim_vars_cpu[Params::V], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  145. std::uniform_real_distribution<Float> gating_var_dist(0.02, 0.05);
  146. for (auto& val : cpu_arr)
  147. val = gating_var_dist(gen);
  148. cudaMemcpy(sim_vars_cpu[Params::M], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  149. for (auto& val : cpu_arr)
  150. val = gating_var_dist(gen);
  151. cudaMemcpy(sim_vars_cpu[Params::H], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  152. for (auto& val : cpu_arr)
  153. val = gating_var_dist(gen);
  154. cudaMemcpy(sim_vars_cpu[Params::N], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  155. for (auto& val : cpu_arr)
  156. val = gating_var_dist(gen);
  157. cudaMemcpy(sim_vars_cpu[Params::P], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  158. for (auto& val : cpu_arr)
  159. val = gating_var_dist(gen);
  160. cudaMemcpy(sim_vars_cpu[Params::Q], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  161. for (auto& val : cpu_arr)
  162. val = gating_var_dist(gen);
  163. cudaMemcpy(sim_vars_cpu[Params::R], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  164. for (auto& val : cpu_arr)
  165. val = gating_var_dist(gen);
  166. cudaMemcpy(sim_vars_cpu[Params::S], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  167. for (auto& val : cpu_arr)
  168. val = gating_var_dist(gen);
  169. cudaMemcpy(sim_vars_cpu[Params::U], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  171. thrust::fill_n(cpu_arr.begin(), N, 0.05);
  172. cudaMemcpy(sim_vars_cpu[Params::IY], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  173. cudaMemcpy(sim_vars_cpu[Params::EY], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  175. thrust::fill_n(cpu_arr.begin(), N, 0.001);
  176. cudaMemcpy(sim_vars_cpu[Params::LOCAL_IOUTPUT], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  177. cudaMemcpy(sim_vars_cpu[Params::LOCAL_IIN], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  178. cudaMemcpy(sim_vars_cpu[Params::LOCAL_ICOUNT], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  179. cudaMemcpy(sim_vars_cpu[Params::LOCAL_EOUTPUT], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  180. cudaMemcpy(sim_vars_cpu[Params::LOCAL_EIN], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  181. cudaMemcpy(sim_vars_cpu[Params::LOCAL_ECOUNT], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  182. cudaMemcpy(sim_vars_cpu[Params::IOUTPUT], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  183. cudaMemcpy(sim_vars_cpu[Params::IIN], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  184. cudaMemcpy(sim_vars_cpu[Params::ICOUNT], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  185. cudaMemcpy(sim_vars_cpu[Params::EOUTPUT], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  186. cudaMemcpy(sim_vars_cpu[Params::EIN], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  187. cudaMemcpy(sim_vars_cpu[Params::ECOUNT], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  189. thrust::fill_n(cpu_arr.begin(), N, 0.0f);
  190. cudaMemcpy(sim_vars_cpu[Params::APPCUR], cpu_arr.data(), N * sizeof(Float), cudaMemcpyHostToDevice);
  191. gpuErrchk(cudaDeviceSynchronize());
  192. sim_vars = sim_vars_cpu;
  193. }
  196. __global__ void init_random(unsigned int seed, int const N, curandState_t* states)
  197. {
  198. for (int idx = blockIdx.x * blockDim.x + threadIdx.x; idx < N; idx += gridDim.x * blockDim.x)
  199. {
  200. curand_init(seed, idx, 0, &states[idx]);
  201. }
  202. }
  204. __global__ void get_random_test(curandState_t* states, int const N, int* numbers)
  205. {
  206. for (int idx = blockIdx.x * blockDim.x + threadIdx.x; idx < N; idx += gridDim.x * blockDim.x)
  207. {
  208. numbers[idx] = curand(&states[idx]) % 100;
  209. }
  210. }
  212. Ephys::Ephys(int const N_, Float const dt_) : N{ N_ }, dt{ dt_ }, random_states(N_)
  213. {
  214. #if USE_FIXED_RANDOM_SEED
  215. unsigned int seed = 0;
  216. #else
  217. unsigned int seed = time(NULL);
  218. #endif
  219. // Setup random seeds.
  220. init_random<<<NBLOCKS, NTHREADS>>>(seed, N, thrust::raw_pointer_cast(random_states.data()));
  221. gpuErrchk( cudaPeekAtLastError() );
  222. gpuErrchk( cudaDeviceSynchronize() );
  223. }
  225. __global__ void Leapfrog_HH(Float** sim_vars, const int N, const Float timestep, curandState_t* random_states)
  226. {
  227. using namespace HH;
  228. Float Vt_d[2] = { -61.5, -61.5};
  229. Float gna_d[2] = { 50, 46};
  230. Float gkd_d[2] = { 4.8, 5.1};
  231. Float gM_d[2] = { 0.13, 0.0};
  232. Float gleak_d[2] = { 0.02, 0.08};
  233. Float tau_max_d[2] = { 1123, 824.5};
  234. Float gT_d[2] = { 0, 0};
  235. Float A_d[2] = { 0.29, 0.22};
  236. for (int idx = blockIdx.x * blockDim.x + threadIdx.x; idx < N; idx += gridDim.x * blockDim.x)
  237. {
  238. Float* V_ptr = sim_vars[Params::V] + idx;
  239. Float* M_ptr = sim_vars[Params::M] + idx;
  240. Float* H_ptr = sim_vars[Params::H] + idx;
  241. Float* N_ptr = sim_vars[Params::N] + idx;
  242. Float* P_ptr = sim_vars[Params::P] + idx;
  243. Float* Q_ptr = sim_vars[Params::Q] + idx;
  244. Float* R_ptr = sim_vars[Params::R] + idx;
  245. Float* S_ptr = sim_vars[Params::S] + idx;
  246. Float* U_ptr = sim_vars[Params::U] + idx;
  247. Float* EY_ptr = sim_vars[Params::EY] + idx;
  248. Float* IY_ptr = sim_vars[Params::IY] + idx;
  250. const Float local_Iin = sim_vars[Params::LOCAL_IIN][idx];
  251. const Float local_Inum_in = sim_vars[Params::LOCAL_ICOUNT][idx];
  252. const Float local_Ein = sim_vars[Params::LOCAL_EIN][idx];
  253. const Float local_Enum_in = sim_vars[Params::LOCAL_ECOUNT][idx];
  254. const Float Iin = sim_vars[Params::IIN][idx];
  255. const Float Inum_in = sim_vars[Params::ICOUNT][idx];
  256. const Float Ein = sim_vars[Params::EIN][idx];
  257. const Float Enum_in = sim_vars[Params::ECOUNT][idx];
  258. const Float appcur = sim_vars[Params::APPCUR][idx];
  259. const Float GAMPA_WHITE = sim_vars[Params::G_AMPA_WHITE][idx];
  260. const Float GAMPA_GREY = sim_vars[Params::G_AMPA_GREY][idx];
  261. const Float GGABA = sim_vars[Params::G_GABA][idx];
  262. const Float ei_info = sim_vars[Params::EI_INFO][idx];
  263. const Float type_info = sim_vars[Params::TYPE_INFO][idx];
  264. const Float V = *V_ptr;
  265. const Float M = *M_ptr;
  266. const Float H = *H_ptr;
  267. const Float N = *N_ptr;
  268. const Float P = *P_ptr;
  269. const Float Q = *Q_ptr;
  270. const Float R = *R_ptr;
  271. const Float S = *S_ptr;
  272. const Float U = *U_ptr;
  273. const Float IY = *IY_ptr;
  274. const Float EY = *EY_ptr;
  275. int neuron_type;
  276. if (type_info < 0.5) {
  277. neuron_type = 0;
  278. }
  279. else {
  280. neuron_type = 1;
  281. }
  283. // Update gating variables first.
  284. Float Vt = Vt_d[neuron_type];
  285. Float aM = -0.32 * (V - Vt - 13) / (exp(-(V - Vt - 13) / 4) - 1);
  286. if ((V - Vt) == 13)
  287. aM = 0.32 * 4;
  288. Float bM = 0.28 * (V - Vt - 40) / (exp((V - Vt - 40) / 5) - 1);
  289. if ((V - Vt) == 40)
  290. bM = 0.28 * 5;
  291. const float new_M = (aM * timestep + (1 - timestep / 2 * (aM + bM)) * M) / (timestep / 2 * (aM + bM) + 1);
  293. Float aH = 0.128 * exp(-(V - Vt - 17) / 18);
  294. Float bH = 4.0 / (1 + exp(-(V - Vt - 40) / 5));
  295. const float new_H = (aH * timestep + (1 - timestep / 2 * (aH + bH)) * H) / (timestep / 2 * (aH + bH) + 1);
  297. Float aN = -0.032 * (V - Vt - 15) / (exp(-(V - Vt - 15) / 5) - 1);
  298. if ((V - Vt) == 15)
  299. aN = 0.032 * 5;
  300. Float bN = 0.5 * exp(-(V - Vt - 10) / 40);
  301. const float new_N = (aN * timestep + (1 - timestep / 2 * (aN + bN)) * N) / (timestep / 2 * (aN + bN) + 1);
  303. Float P_inf = 1 / (1 + exp(-(V + 35) / 10));
  304. Float tau_P = tau_max_d[neuron_type] / (3.3 * exp((V + 35) / 20) + exp(-(V + 35) / 20));
  305. const float new_P = 2 * timestep / (2 * tau_P + timestep) * P_inf + (2 * tau_P - timestep) / (2 * tau_P + timestep) * P;
  307. Float aQ = 0.055 * (-27 - V) / (exp((-27 - V) / 3.8) - 1);
  308. if (V == -27)
  309. aQ = 0.055 * 3.8;
  310. Float bQ = 0.94 * exp((-75 - V) / 17);
  311. const float new_Q = (aQ * timestep + (1 - timestep / 2 * (aQ + bQ)) * Q) / (timestep / 2 * (aQ + bQ) + 1);
  313. Float aR = 0.000457 * exp((-13 - V) / 50);
  314. Float bR = 0.0065 / (exp((-15 - V) / 28) + 1);
  315. const float new_R = (aR * timestep + (1 - timestep / 2 * (aR + bR)) * R) / (timestep / 2 * (aR + bR) + 1);
  318. Float S_inf; Float U_inf; Float tau_U; Float tau_S;
  321. if (neuron_type == 2) {
  322. S_inf = 1 / (1 + exp(-(V + 59) / 6.2));
  323. U_inf = 1 / (1 + exp((V + 83) / 4));
  324. tau_S = 0.13 + 0.22 / (exp(-(V + 132) / 16.7) + exp((V + 16.8) / 18.2));
  325. tau_U = 8.2 + (56.6 + 0.27 * exp((V + 115.2) / 5)) / (1 + exp((V + 86) / 3.2));
  326. }
  327. else {
  328. S_inf = 1 / (1 + exp(-(V + 52) / 7.4));
  329. U_inf = 1 / (1 + exp((V + 80) / 5));
  330. tau_S = 1 + 0.33 / (exp(-(V + 102) / 15) + exp((V + 27) / 10));
  331. tau_U = 22.7 + 0.27 / (exp(-(V + 407) / 50) + exp((V + 48) / 4));
  332. }
  335. const float new_U = 2 * timestep / (2 * tau_U + timestep) * U_inf + (2 * tau_U - timestep) / (2 * tau_U + timestep) * U;
  336. const float new_S = 2 * timestep / (2 * tau_S + timestep) * S_inf + (2 * tau_S - timestep) / (2 * tau_S + timestep) * S;
  338. const Float Iin_avg = Iin + local_Iin; //(Inum_in + local_Inum_in) > 0 ? (Iin + local_Iin) / (Inum_in + local_Inum_in) : 0;//
  339. const Float new_IY = (gabaR * Iin_avg * timestep + (1 - timestep / 2 * (gabaR * Iin_avg + gabaD)) * IY) / (timestep / 2 * (gabaR * Iin_avg + gabaD) + 1);
  340. const Float Ein_avg = Ein + local_Ein; //(Enum_in + local_Enum_in) > 0 ? (Ein + local_Ein) / (Enum_in + local_Enum_in) : 0; //
  341. const Float new_EY = (ampaR * Ein_avg * timestep + (1 - timestep / 2 * (ampaR * Ein_avg + ampaD)) * EY) / (timestep / 2 * (ampaR * Ein_avg + ampaD) + 1);
  343. // Use updated gating variables to update voltage.
  344. const Float leak_term = gleak_d[neuron_type];
  345. Float g_na = gna_d[neuron_type];
  346. const Float Na_term = g_na * new_M * new_M * new_M * new_H;
  347. Float g_kd = gkd_d[neuron_type];
  348. const Float Kd_term = g_kd * powf(new_N, 4);
  349. Float g_M = gM_d[neuron_type];
  351. const Float IM_term = g_M* new_P;
  352. const Float IL_term = 0;// gca* new_Q* new_Q* new_R;
  353. const Float IT_term = gT_d[neuron_type] * new_S* new_S* new_U;
  354. Float g_GABA = GGABA;
  356. if (type_info < 1.5)
  357. g_GABA = 0.1;
  359. const Float global_gaba_term = g_GABA * new_IY;
  362. const Float global_ampa_term = GAMPA_WHITE * new_EY;
  366. const Float G = leak_term + Na_term + Kd_term + IM_term + IL_term + IT_term + global_gaba_term + global_ampa_term; //local_gaba_term + local_ampa_term +
  367. const Float E = leak_term * el + Na_term * ena + Kd_term * ek + IM_term * ek + (IL_term +IT_term) * eca + + global_gaba_term * e_gaba + global_ampa_term * e_ampa; //+ local_gaba_term * e_gaba + local_ampa_term * e_ampa
  369. // Add random white noise to V.
  370. Float white_noise = white_noise_intensity * curand_normal(random_states + idx) + white_noise_mean;
  372. // Copy back to sim_vars
  373. *V_ptr = ((white_noise + appcur) * timestep + E * timestep + (1 - timestep / 2 * G)*V) / ( C * (1 + timestep / 2 * G));
  374. *M_ptr = new_M;
  375. *H_ptr = new_H;
  376. *N_ptr = new_N;
  377. *P_ptr = new_P;
  378. *Q_ptr = new_Q;
  379. *R_ptr = new_R;
  380. *S_ptr = new_S;
  381. *U_ptr = new_U;
  382. *IY_ptr = new_IY;
  383. *EY_ptr = new_EY;
  385. // Calculate amount of neurotransmitter outputted AFTER update
  386. if (ei_info < 0.5) {
  387. sim_vars[Params::EOUTPUT][idx] = sim_vars[Params::LOCAL_EOUTPUT][idx] = 1.0 / (1.0 + exp(-(*V_ptr - Vtc) / Kp));
  388. sim_vars[Params::IOUTPUT][idx] = sim_vars[Params::LOCAL_IOUTPUT][idx] = 0;
  389. }
  390. else {
  391. sim_vars[Params::IOUTPUT][idx] = sim_vars[Params::LOCAL_IOUTPUT][idx] = 1.0 / (1.0 + exp(-(*V_ptr - Vtc) / Kp));
  392. sim_vars[Params::EOUTPUT][idx] = sim_vars[Params::LOCAL_EOUTPUT][idx] = 0;
  393. }
  395. }
  396. }
  398. void HHEphys::SimulateEphys()
  399. {
  400. Leapfrog_HH<<<NBLOCKS, NTHREADS>>>(thrust::raw_pointer_cast(sim_vars.data()), N, dt, thrust::raw_pointer_cast(random_states.data()));
  401. gpuErrchk( cudaPeekAtLastError() );
  402. gpuErrchk( cudaDeviceSynchronize() );
  403. }

ephys.cu, under CC-BY-4.0 · at the source

Overview

Authors: Guanhua Sun1, James Hazelden1, Ruby Kim1, Daniel B Forger1,2
ORCID iDs: Guanhua Sun, James Hazelden, Ruby Kim, Daniel B Forger
  1. Department of Mathematics, University of Michigan Ann Arbor United States
  2. Gilbert S. Omenn Department of Computational Medicine and Bioinformatics, University of Michigan Ann Arbor United States
Institutions: University of Michigan (United States)
Journal: eLife, volume 14, article RP108208
Dates: published online 4 September 2026
Type: Research article · Language: English
License: CC BY
Identifiers: DOI 10.7554/elife.108208 · PMID 42696350 · PMCID PMC13544912 · OpenAlex W4414145370
Open access: gold, a free copy (OpenAlex)
Status: code verified
Categories: mouse (organism), computational (subfield)
Methods: Spectral & time-frequency, Machine learning
Keywords: traveling waves, computational modeling, network models, Mouse
MeSH: Brain Waves*, Cerebral Cortex*, Models, Neurological*, Nerve Net*, Neurons*, Animals, Mice (* major topic)
Journal subjects: Neuroscience
Topic: Neural dynamics and brain function (Cognitive Neuroscience, Neuroscience), according to OpenAlex
Funding: Army Research Office (W911NF-22-1-0223); Human Frontier Science Program (10.52044/hfsp.rgp00192018.pc.gr.166630); National Science Foundation (DMS 2052499)
Citations: cited by 1 paper (Europe PMC); 66 references in the paper

Abstract

Traveling waves are ubiquitous in neuronal systems across different spatial scales. While microscopic and mesoscopic waves are relatively well studied, the emergence of macroscopic traveling waves remains less understood. Here, by modeling the mouse cortex using spatial transcriptomic and connectivity data, we show that realistic cortical connectivity can generate a significantly higher level of macroscopic traveling waves than artificial local and uniform connectivity across multiple oscillation frequency bands, with the strongest advantage appearing in the theta, alpha, and beta frequency bands. By probing the model in different dynamic regimes, we find that macroscopic wave activity depends on both network connectivity and excitatory coupling strength, with a non-monotonic dependence on coupling. Together, our work shows how flexible macroscopic traveling waves can emerge in the mouse cortex and offers a computational framework to further study traveling waves in the mouse brain at the single-cell level.

Reproduced under the paper's license (CC BY), from the paper cited above.

Repository

Its files are read in the Code ↔ Paper reader above, with 1 match between paragraphs and lines of code.

figshare 29410397

License: CC-BY-4.0
State: the link answers, verified on 26 September 2026
Evidence: files inventoried
Size: 13 files
Software Heritage: not checked
Found in: “Data availability”
Not found: README, license file, CITATION.cff, environment file, tests, continuous integration, documentation
Tools: AllenSDK (2 files), h5py (2 files), NumPy (2 files), pandas (2 files)
Availability: 1 check, the latest on 26 September 2026: the link answers (HTTP 200)
  • 26 September 2026: the link answers (HTTP 200)
16 files
At the source: doi.org/10.6084/m9.figshare.29410397

The paper's code and data availability statement is in the Data section.

Tracing map

Proposed by the machine: these links were found in the paper and verified at the source, without human review. The map will receive a Zenodo DOI once one of the paper's authors has validated it with their ORCID.

What the map holds:

  • 1 repository of the authors' code, each at its verified commit, with its license and how the link was found in the paper;
  • 14 scripts, each with its path and the digest of its content;
  • 1 match between paragraphs of the paper and lines of the code (method lexical-v1);
  • neither the text of the paper nor the code itself.

Its JSON (tracing-map.json) is deposited on Zenodo with its DOI once the map is validated.

Data

No dataset and no data link were found in the paper.

Data availability

Raw data related to this study, including the neuronal and connectivity data, along with the code and simulation data related to the figures, have been uploaded to the following repository: https://doi.org/10.6084/m9.figshare.29410397.

The following dataset was generated:

SunG 2026Realistic coupling enables flexible macroscopic traveling waves in the mouse cortexfigshare10.6084/m9.figshare.29410397.v2PMC1354491242696350

Reproduced under the paper's license (CC BY), from the paper cited above.

Versions

The history of this record: each version stored by the harvester or made by a correction of its authors or of the maintainers of its code, and what changed in its facts. The texts of the paper (its abstract, its availability statements) are not part of it; versions that changed only those are not listed.

Version 1, 27 September 2026: the first record

Recorded: type, language, journal, volume, pages, dates, 4 authors, 4 keywords, 7 MeSH terms, 3 funders, 65 references.

Cite

This paper

Sun, G., Hazelden, J., Kim, R., & Forger, D. B. (2026). Realistic coupling enables flexible macroscopic traveling waves in the mouse cortex. eLife, 14, RP108208. https://doi.org/10.7554/elife.108208

BibTeX

@article{sun2026realistic,
author = {Sun, Guanhua and Hazelden, James and Kim, Ruby and Forger, Daniel B},
title = {{Realistic coupling enables flexible macroscopic traveling waves in the mouse cortex}},
journal = {eLife},
year = {2026},
month = sep,
volume = {14},
pages = {RP108208},
publisher = {eLife Sciences Publications, Ltd},
issn = {2050-084X},
doi = {10.7554/elife.108208},
url = {https://doi.org/10.7554/elife.108208},
pmid = {42696350},
pmcid = {PMC13544912}
}

RIS

TY - JOUR
AU - Sun, Guanhua
AU - Hazelden, James
AU - Kim, Ruby
AU - Forger, Daniel B
TI - Realistic coupling enables flexible macroscopic traveling waves in the mouse cortex
T2 - eLife
J2 - Elife
PY - 2026
DA - 2026/09/04
VL - 14
SP - RP108208
SN - 2050-084X
PB - eLife Sciences Publications, Ltd
DO - 10.7554/elife.108208
UR - https://doi.org/10.7554/elife.108208
LA - en
ER -

CSL-JSON

{
"id": "10.7554/elife.108208",
"type": "article-journal",
"title": "Realistic coupling enables flexible macroscopic traveling waves in the mouse cortex",
"container-title": "eLife",
"author": [
{
"family": "Sun",
"given": "Guanhua"
},
{
"family": "Hazelden",
"given": "James"
},
{
"family": "Kim",
"given": "Ruby"
},
{
"family": "Forger",
"given": "Daniel B"
}
],
"container-title-short": "Elife",
"volume": "14",
"page": "RP108208",
"DOI": "10.7554/elife.108208",
"PMID": "42696350",
"PMCID": "PMC13544912",
"ISSN": "2050-084X",
"publisher": "eLife Sciences Publications, Ltd",
"URL": "https://doi.org/10.7554/elife.108208",
"language": "en",
"issued": {
"date-parts": [
[
2026,
9,
4
]
]
}
}

The tracing map gets a citation of its own once an author has validated it and it has a DOI.

Similar papers

The papers with a page that share the most with this one: the tools found in their code, their categories, datasets, cited references and authors, the rarest counting most.

[1] doi:10.1038/s41467-026-71386-z [code]
Planar, spiral, and concentric traveling waves distinguish behavioral states in human memory.
Journal: Nature communications
In common: pandas, NumPy, 17 references
[2] doi:10.3389/fncom.2026.1844662
Quantifying cortex-wide traveling brain waves of complex patterns with a graph-based algorithm.
Journal: Frontiers in computational neuroscience
In common: computational, 13 references
[3] doi:10.1016/j.isci.2026.116728 [code]
Awake cortex stabilizes traveling waves for global and reliable information routing.
Journal: iScience
In common: NumPy, 9 references
[4] doi:10.7554/elife.106753 [code]
Traveling waves across scales: Different mechanisms but same canonical computation?
Journal: n/a
In common: computational, 8 references
[5] doi:10.1073/pnas.2527296123
Traveling-wave transcranial alternating current stimulation (twtACS) causally links neural timing to cognitive function.
Journal: Proceedings of the National Academy of Sciences of the United States of America
In common: 8 references
[6] doi:10.7554/elife.100674 [code]
The dominance of large-scale phase dynamics in human cortex, from delta to gamma.
Journal: eLife
In common: NumPy, 7 references
[7] doi:10.1162/imag.a.1310 [code]
Experimental quality control induces changes in Allen mouse brain connectomes.
Journal: Imaging neuroscience (Cambridge, Mass.)
In common: AllenSDK, pandas, NumPy, mouse, 4 references
[8] doi:10.1038/s41593-026-02253-9 [code]
Genoarchitecture and input-output organization of the mouse basal ganglia and thalamic parafascicular nucleus.
Journal: Nature neuroscience
In common: AllenSDK, h5py, pandas, 1 other tool, mouse, 3 references
[9] doi:10.1371/journal.pcbi.1013007 [code]
Traveling waves in the human visual cortex: An MEG-EEG model-based approach
Journal: n/a
In common: NumPy, 6 references
[10] doi:10.7554/elife.95562 [code]
Spatially targeted inhibitory rhythms differentially affect neuronal integration.
Journal: eLife
In common: h5py, pandas, NumPy, 4 references

Contribute

The authors of this paper can claim it, correct its record and validate its tracing map, and the maintainers of its code (its owner, or a public member of its organization) correct what it says of their repository; anyone signed in can ask for its removal. Every request goes to OSCR's own machine, which answers it; your account page follows them.

Sign in with ORCID to claim this paper as one of its authors, correct its record or validate its tracing map: when the paper's metadata lists your ORCID iD, you are recognized at once. Maintainers of its code: sign in with GitHub, then claim the repository on your account page.

Request its removal

To ask OSCR to remove this record, the copies of its authors' scripts or its tracing map, use the removal request page: signed in, you say who you are, what to remove and why, then review and confirm the request. Published rules decide every request (how).

Discussion, reproductions, activity

Discussion: questions and error reports about this paper and its code, from signed-in readers and its authors. It opens with sign-in.

Reproductions: reports from readers who ran the authors' code: what they reproduced, with which environment, commit and data. It opens with sign-in.

Activity: what happens around this paper: new versions of its record, its map's validation, discussions and reproductions. It opens with sign-in.