Skip to content
This repository was archived by the owner on Jun 8, 2023. It is now read-only.

Commit 42bb3c3

Browse files
authored
Support fast_imem on GPU. (#64)
1 parent 92d62cc commit 42bb3c3

File tree

1 file changed

+14
-2
lines changed

1 file changed

+14
-2
lines changed

src/mod2c_core/noccout.c

Lines changed: 14 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -153,7 +153,19 @@ static void rhs_d_pnt_race(const char* r, const char* d) {
153153
\n #pragma acc atomic update\
154154
\n _vec_rhs[_nd_idx] %s _rhs;\
155155
\n #pragma acc atomic update\
156-
\n _vec_d[_nd_idx] %s _g;\
156+
\n _vec_d[_nd_idx] %s _g;", r, d);
157+
P(buf);
158+
if(electrode_current) {
159+
sprintf(buf, "\
160+
\n if (_nt->nrn_fast_imem) {\
161+
\n #pragma acc atomic update\
162+
\n _nt->nrn_fast_imem->nrn_sav_rhs[_nd_idx] %s _rhs;\
163+
\n #pragma acc atomic update\
164+
\n _nt->nrn_fast_imem->nrn_sav_d[_nd_idx] %s _g;\
165+
\n }", r, d);
166+
P(buf);
167+
}
168+
sprintf(buf, "\
157169
\n } else {\
158170
\n _vec_shadow_rhs[_iml] = _rhs;\
159171
\n _vec_shadow_d[_iml] = _g;\
@@ -177,7 +189,7 @@ static void rhs_d_pnt_race(const char* r, const char* d) {
177189
\n _vec_d[_nd_idx] %s _vec_shadow_d[_iml];\
178190
%s\
179191
\n#endif\
180-
\n", r, d, r, d, r, d, print_fast_imem_code());
192+
\n", r, d, r, d, print_fast_imem_code());
181193
P(buf);
182194
}
183195

0 commit comments

Comments
 (0)