Neko 1.99.5
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
gs_utofu_aux.c
Go to the documentation of this file.
1/*
2 Copyright (c) 2026, The Neko Authors
3 All rights reserved.
4
5 Redistribution and use in source and binary forms, with or without
6 modification, are permitted provided that the following conditions
7 are met:
8
9 * Redistributions of source code must retain the above copyright
10 notice, this list of conditions and the following disclaimer.
11
12 * Redistributions in binary form must reproduce the above
13 copyright notice, this list of conditions and the following
14 disclaimer in the documentation and/or other materials provided
15 with the distribution.
16
17 * Neither the name of the authors nor the names of its
18 contributors may be used to endorse or promote products derived
19 from this software without specific prior written permission.
20
21 THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
22 "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
23 LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
24 FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
25 COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
26 INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
27 BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
28 LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
29 CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
30 LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
31 ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
32 POSSIBILITY OF SUCH DAMAGE.
33*/
34
70#include <stdlib.h>
71#include <stdint.h>
72#include <utofu.h>
73
74/* Put flags: notice on the local TCQ (so we can reclaim the send buffer)
75 * and on the remote MRQ (so the receiver learns of arrival). Cache
76 * injection (added at init unless NEKO_GS_UTOFU_CACHE_INJECT=0) drops the
77 * received slab straight into the peer's cache -- a win if it is consumed
78 * immediately, a loss if it pollutes cache, hence the runtime toggle.
79 *
80 * UTOFU_ONESIDED_FLAG_STRONG_ORDER (which jacobi2d sets) is deliberately
81 * NOT used: it is only required when arrival is detected by watchdog-
82 * polling the payload's last word, which needs the data to land in address
83 * order. The RMT_PUT MRQ notice used here is raised only after the whole
84 * put has been written at the target, so it already carries that
85 * guarantee, and strong ordering would just restrict the NIC's packet-
86 * level freedom on multi-packet slabs. Revisit if the receive side ever
87 * moves to jacobi2d-style payload watchdogs. */
88#define GS_UTOFU_PUT_FLAGS_BASE \
89 (UTOFU_ONESIDED_FLAG_TCQ_NOTICE | UTOFU_ONESIDED_FLAG_REMOTE_MRQ_NOTICE)
90
91/* --- process-global endpoint state, created once ------------------------ */
92static int gs_utofu_ready = 0;
93static size_t gs_utofu_nvcq = 0; /* number of injection VCQs */
94static size_t gs_utofu_ntni = 0; /* TNIs used for injection */
95static utofu_vcq_hdl_t *gs_utofu_vcq = NULL; /* [nvcq] injection VCQs */
96static utofu_tni_id_t *gs_utofu_alltni = NULL; /* [nalltni] ALL 1-sided TNIs;
97 * first ntni = injection set */
98static size_t gs_utofu_nalltni = 0;
99static int gs_utofu_recv_rr = 0; /* round-robins recv-VCQ TNIs */
103static size_t gs_utofu_nrvcq_def = 1; /* recv VCQs per ctx (req.) */
104
105/* Longs per counter slot: one A64FX L1/L2 cache line (256 bytes), so the
106 * per-VCQ completion counters never false-share between threads. */
107#define GS_UTOFU_PEND_STRIDE 32
108
109/* --- per-instance context (one per gs_utofu_t) -------------------------- */
110typedef struct {
111 utofu_stadd_t *send_stadd; /* [nvcq] send_buf registered on each inj VCQ */
112 /* This instance's OWN receive VCQs, spread over the TNIs. Each receive
113 * peer is bound to one of them at init (the receiver advertises that
114 * VCQ's id/STADD to the peer), so incoming halo traffic is processed by
115 * several TNIs instead of funnelling through one -- the receive-side
116 * counterpart of the per-thread injection pool. All are polled by the
117 * master; the edata tags are per-instance-unique, so which MRQ a notice
118 * appears on carries no protocol meaning. */
119 size_t nrvcq; /* receive VCQs actually created (>= 1) */
120 utofu_vcq_hdl_t *recv_vcq; /* [nrvcq] */
121 utofu_stadd_t *recv_stadd; /* [nrvcq] recv_buf registered on each */
122 /* Outstanding (un-reaped) put completions, one padded slot per
123 * (injection VCQ, send-buffer half). The put's cbdata points at its
124 * slot, so draining a VCQ's TCQ decrements the right counter -- letting
125 * nbsend wait only on the half it is about to reuse. The phase barriers
126 * in gs_nbsend_utofu guarantee each VCQ (and thus each slot) has a
127 * single toucher at any time, so no atomics are needed. */
128 long *send_pending; /* [2 * nvcq * GS_UTOFU_PEND_STRIDE] */
129 /* notices observed early that belong to the *other* parity (next round) */
134
135static long *gs_utofu_pend_slot(gs_utofu_ctx_t *ctx, size_t k, int half)
136{
137 return &ctx->send_pending[(2 * k + (size_t) half) * GS_UTOFU_PEND_STRIDE];
138}
139
140/* Drain every completion currently sitting on VCQ k's transmit queue,
141 * decrementing the per-(instance, VCQ, half) counter each put carried as
142 * cbdata. Must only be called by VCQ k's single toucher of the current
143 * phase. */
144static void gs_utofu_drain_tcq(size_t k)
145{
146 void *cbdata;
148 if (cbdata != NULL) (*(long *) cbdata)--;
149 }
150}
151
153{
154 if (ctx->stash_n == ctx->stash_cap) {
155 int cap = ctx->stash_cap ? 2 * ctx->stash_cap : 16;
156 ctx->stash = realloc(ctx->stash, cap * sizeof(*ctx->stash));
157 ctx->stash_cap = cap;
158 }
159 ctx->stash[ctx->stash_n++] = edata;
160}
161
181 int *nrvcq_out, uint64_t *edata_max, int *ierr)
182{
183 *ierr = 0;
184
185 if (!gs_utofu_ready) {
187 size_t num_tnis = 0;
188
190 || num_tnis == 0) {
191 *ierr = 1;
192 return;
193 }
194
195 size_t want_tni = (ntni_req > 0) ? (size_t) ntni_req : 1;
200 for (size_t k = 0; k < num_tnis; k++)
202
203 /* One injection VCQ per thread, capped by NEKO_GS_UTOFU_NVCQ. The VCQ
204 * budget is shared with the per-instance receive VCQs created later
205 * (one scalar + one vector per gs instance) -- on Fugaku a TNI grants
206 * us ~8 VCQs, so a 12-thread pool can eat a whole TNI; the cap lets a
207 * tight budget be rebalanced without rebuilding. THREAD_SAFE is what
208 * makes multithreaded operation defined at all: without it uTofu takes
209 * the unlocked direct-descriptor injection path and any cross-thread
210 * touch corrupts the TOQ (the "TOQ Direct Descriptor Exception" of the
211 * first multithreaded attempt). With one thread per VCQ the internal
212 * lock is uncontended. */
213 size_t want_vcq = (nthreads > 0) ? (size_t) nthreads : 1;
214 const char *nv = getenv("NEKO_GS_UTOFU_NVCQ");
215 if (nv != NULL) {
216 long cap = strtol(nv, NULL, 10);
217 if (cap > 0 && (size_t) cap < want_vcq)
218 want_vcq = (size_t) cap;
219 }
220
221 /* Preferred TNI first, then the rest of the injection set: one
222 * exhausted TNI must not end the pool while others still have room.
223 * If every TNI in the set is out of VCQs, keep what was granted:
224 * threads beyond the count simply do not inject (see
225 * gs_utofu_post_puts), and the modulo peer partition shrinks with it. */
227 while (gs_utofu_nvcq < want_vcq) {
228 int created = 0;
229 for (size_t s = 0; s < want_tni && !created; s++) {
233 created = 1;
234 }
235 if (!created)
236 break;
238 }
239 if (gs_utofu_nvcq == 0) {
240 *ierr = 2;
241 return;
242 }
243
244 /* Start the receive-VCQ round-robin just past the injection TNIs, so
245 * receive queues prefer interfaces the pool has not loaded. */
247
248 /* Receive VCQs per instance path: default one per TNI, so an
249 * instance's incoming halo traffic is processed by every interface
250 * instead of funnelling through one. NEKO_GS_UTOFU_NRVCQ overrides
251 * (=1 restores the single-receive-VCQ behaviour). The budget is
252 * shared with the injection pool and across instances/paths;
253 * ctx_create degrades gracefully to fewer when it runs dry. */
255 const char *nr = getenv("NEKO_GS_UTOFU_NRVCQ");
256 if (nr != NULL) {
257 long v = strtol(nr, NULL, 10);
258 if (v > 0 && v <= 64) gs_utofu_nrvcq_def = (size_t) v;
259 }
260
261 /* Largest edata the hardware will carry. VERIFY: field name and whether
262 * the capability is expressed in bytes (assumed here). */
263 struct utofu_onesided_caps *caps = NULL;
265 && caps != NULL) {
266 size_t eb = caps->max_edata_size;
267 /* Capped to a positive signed-64 value: the Fortran side receives this
268 * as integer(c_int64_t) and compares against it. */
270 else if (eb > 0) gs_utofu_edata_max = (1ULL << (8 * eb)) - 1ULL;
271 }
272
273 /* Cache injection on by default; NEKO_GS_UTOFU_CACHE_INJECT=0 disables. */
274 const char *ci = getenv("NEKO_GS_UTOFU_CACHE_INJECT");
276 if (ci == NULL || ci[0] != '0')
278
279 /* A/B switch: serialise all injection onto thread 0 (still spread over
280 * every VCQ, which is legal now that they are THREAD_SAFE) to isolate
281 * the effect of parallel injection against the same transport. */
282 const char *mi = getenv("NEKO_GS_UTOFU_MASTER_INJECT");
283 gs_utofu_master_inject = (mi != NULL && mi[0] != '0');
284
285 free(tnis);
286 gs_utofu_ready = 1;
287 }
288
293}
294
310void gs_utofu_ctx_create(void *send_buf, size_t send_bytes,
311 void *recv_buf, size_t recv_bytes,
312 void **ctx_out, int *nrvcq_out,
313 uint64_t *recv_vcq_id, uint64_t *recv_stadd,
314 int *ierr)
315{
316 *ierr = 0;
317 gs_utofu_ctx_t *ctx = calloc(1, sizeof(*ctx));
318 ctx->send_stadd = malloc(gs_utofu_nvcq * sizeof(*ctx->send_stadd));
320 sizeof(*ctx->send_pending));
321 ctx->recv_vcq = malloc(gs_utofu_nrvcq_def * sizeof(*ctx->recv_vcq));
322 ctx->recv_stadd = malloc(gs_utofu_nrvcq_def * sizeof(*ctx->recv_stadd));
323
324 /* Receive VCQs may sit on ANY one-sided TNI: incoming puts are routed
325 * by the advertised VCQ id, so spreading the receive processing over all
326 * interfaces is free (NEKO_GS_UTOFU_NTNI deliberately confines only the
327 * injection side). Round-robin over the full set, sweeping past
328 * exhausted TNIs -- the injection pool can eat a whole TNI's VCQ budget
329 * (~8 per TNI on Fugaku). Stop early when the budget runs dry; at least
330 * one must succeed, since a private MRQ per instance is a correctness
331 * requirement. They stay non-THREAD_SAFE (flags 0): they are strictly
332 * master-polled, and skipping the internal lock keeps the MRQ sweep --
333 * the receive hot path -- cheap. Never touch them off-master. */
334 while (ctx->nrvcq < gs_utofu_nrvcq_def) {
335 int created = 0;
336 for (size_t s = 0; s < gs_utofu_nalltni && !created; s++) {
339 if (utofu_create_vcq(rtni, 0, &ctx->recv_vcq[ctx->nrvcq])
340 == UTOFU_SUCCESS)
341 created = 1;
342 }
344 if (!created)
345 break;
346 ctx->nrvcq++;
347 }
348 if (ctx->nrvcq == 0) {
349 *ierr = 1;
350 return;
351 }
352
353 for (size_t k = 0; k < gs_utofu_nvcq; k++) {
354 if (utofu_reg_mem(gs_utofu_vcq[k], send_buf, send_bytes, 0,
355 &ctx->send_stadd[k]) != UTOFU_SUCCESS) {
356 *ierr = 2;
357 return;
358 }
359 }
360 for (size_t r = 0; r < ctx->nrvcq; r++) {
361 if (utofu_reg_mem(ctx->recv_vcq[r], recv_buf, recv_bytes, 0,
362 &ctx->recv_stadd[r]) != UTOFU_SUCCESS) {
363 *ierr = 3;
364 return;
365 }
366 }
367
368 for (size_t r = 0; r < ctx->nrvcq; r++) {
370 if (utofu_query_vcq_id(ctx->recv_vcq[r], &vid) != UTOFU_SUCCESS) {
371 *ierr = 4;
372 return;
373 }
374 recv_vcq_id[r] = (uint64_t) vid;
375 recv_stadd[r] = (uint64_t) ctx->recv_stadd[r];
376 }
377
378 *nrvcq_out = (int) ctx->nrvcq;
379 *ctx_out = ctx;
380}
381
382void gs_utofu_ctx_free(void *vctx, int *ierr)
383{
384 *ierr = 0;
385 gs_utofu_ctx_t *ctx = vctx;
386 if (ctx == NULL) return;
387
388 /* Make sure no put is still reading send_buf before we deregister it.
389 * Called from a serial region, so one thread draining every VCQ is safe
390 * (the VCQs are THREAD_SAFE, and there is no concurrent toucher). */
391 for (int h = 0; h < 2; h++)
392 for (size_t k = 0; k < gs_utofu_nvcq; k++)
393 while (*gs_utofu_pend_slot(ctx, k, h) > 0)
395
396 for (size_t k = 0; k < gs_utofu_nvcq; k++)
398 for (size_t r = 0; r < ctx->nrvcq; r++) {
399 utofu_dereg_mem(ctx->recv_vcq[r], ctx->recv_stadd[r], 0);
400 utofu_free_vcq(ctx->recv_vcq[r]);
401 }
402
403 free(ctx->send_stadd);
404 free(ctx->recv_vcq);
405 free(ctx->recv_stadd);
406 free(ctx->send_pending);
407 free(ctx->stash);
408 free(ctx);
409}
410
411/* Post peer i's put on injection VCQ k, charging its completion to the
412 * (k, parity) counter slot so the next reuse of this send-buffer half waits
413 * on exactly its own puts. The caller must be VCQ k's single toucher of the
414 * current phase (the BUSY-retry path drains k's TCQ). */
415static int gs_utofu_put_one(gs_utofu_ctx_t *ctx, size_t k, int i,
416 const int64_t *send_off, const int64_t *send_len,
417 const uint64_t *rmt_vcq_id,
418 const uint64_t *rmt_stadd,
419 const int64_t *dst_off0, const int64_t *dst_off1,
420 const int *dst_tag, int parity, int64_t send_base,
421 int elem_size)
422{
423 const int64_t doff = parity ? dst_off1[i] : dst_off0[i];
424 long *pend = gs_utofu_pend_slot(ctx, k, parity);
425
427 + (utofu_stadd_t)((send_base + send_off[i]) * elem_size);
428 utofu_stadd_t rmt = (utofu_stadd_t) rmt_stadd[i]
429 + (utofu_stadd_t)(doff * elem_size);
430 size_t len = (size_t)(send_len[i] * elem_size);
431 uint64_t edata = (uint64_t) dst_tag[i] * 2u + (uint64_t) parity;
432
433 int rc;
434 do {
435 rc = utofu_put(gs_utofu_vcq[k], (utofu_vcq_id_t) rmt_vcq_id[i],
437 /* TOQ full (reported as UTOFU_ERR_BUSY): reclaim completions, retry. */
438 if (rc == UTOFU_ERR_BUSY)
440 } while (rc == UTOFU_ERR_BUSY);
441
442 if (rc == UTOFU_SUCCESS)
443 (*pend)++;
444 return rc;
445}
446
463void gs_utofu_post_puts(void *vctx, int tid, int nteam, int nsend,
464 const int64_t *send_off, const int64_t *send_len,
465 const uint64_t *rmt_vcq_id, const uint64_t *rmt_stadd,
466 const int64_t *dst_off0, const int64_t *dst_off1,
467 const int *dst_tag, int parity, int64_t send_base,
468 int elem_size, int *ierr)
469{
470 *ierr = 0;
471 gs_utofu_ctx_t *ctx = vctx;
472
474 if (tid != 0) return;
475 for (int i = 0; i < nsend; i++) {
476 size_t k = (size_t)(i % (int) gs_utofu_nvcq);
477 if (gs_utofu_put_one(ctx, k, i, send_off, send_len, rmt_vcq_id,
478 rmt_stadd, dst_off0, dst_off1, dst_tag, parity,
479 send_base, elem_size) != UTOFU_SUCCESS) {
480 *ierr = 1;
481 return;
482 }
483 }
484 return;
485 }
486
488 if (tid >= npost) return;
489
490 for (int i = tid; i < nsend; i += npost) {
491 if (gs_utofu_put_one(ctx, (size_t) tid, i, send_off, send_len,
492 rmt_vcq_id, rmt_stadd, dst_off0, dst_off1, dst_tag,
493 parity, send_base, elem_size) != UTOFU_SUCCESS) {
494 *ierr = 1;
495 return;
496 }
497 }
498}
499
509void gs_utofu_poll_recv(void *vctx, int parity, int cap,
510 int *ncompleted, int *out_idx, int *ierr)
511{
512 *ierr = 0;
513 gs_utofu_ctx_t *ctx = vctx;
514 int n = 0;
515
516 /* Replay stashed notices that now match the current parity. */
517 int w = 0;
518 for (int s = 0; s < ctx->stash_n; s++) {
519 uint64_t ed = ctx->stash[s];
520 if ((int)(ed & 1u) == parity && n < cap)
521 out_idx[n++] = (int)(ed >> 1) + 1;
522 else
523 ctx->stash[w++] = ed;
524 }
525 ctx->stash_n = w;
526
527 /* Drain currently-available arrival notices from every one of this
528 * instance's MRQs. Each peer's notices always land on its bound VCQ, but
529 * the tags are instance-unique, so the queues can be drained in any
530 * order into the one result list / stash. */
532 for (size_t r = 0; r < ctx->nrvcq && n < cap; r++) {
533 while (n < cap) {
534 int rc = utofu_poll_mrq(ctx->recv_vcq[r], 0, &notice);
535 if (rc == UTOFU_ERR_NOT_FOUND) break;
536 if (rc != UTOFU_SUCCESS) {
537 *ierr = 1;
538 return;
539 }
540 /* VERIFY: RMT_PUT is the receiver-side notice for an incoming put. */
541 if (notice.notice_type != UTOFU_MRQ_TYPE_RMT_PUT) continue;
542
543 uint64_t ed = notice.edata;
544 if ((int)(ed & 1u) == parity)
545 out_idx[n++] = (int)(ed >> 1) + 1;
546 else
548 }
549 }
550
551 *ncompleted = n;
552}
553
562void gs_utofu_drain_half(void *vctx, int half, int tid, int nteam, int *ierr)
563{
564 *ierr = 0;
565 gs_utofu_ctx_t *ctx = vctx;
566 for (size_t k = (size_t) tid; k < gs_utofu_nvcq; k += (size_t) nteam)
567 while (*gs_utofu_pend_slot(ctx, k, half) > 0)
569}
__global__ void ale_add_kinematics_kernel(const int n, T *__restrict__ wx, T *__restrict__ wy, T *__restrict__ wz, const T *__restrict__ x_ref, const T *__restrict__ y_ref, const T *__restrict__ z_ref, const T *__restrict__ phi, const T *__restrict__ x, const T *__restrict__ y, const T *__restrict__ z, const kinematics_params_t kin_params)
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w
const int i
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v
static int gs_utofu_ready
static void gs_utofu_stash_push(gs_utofu_ctx_t *ctx, uint64_t edata)
static int gs_utofu_master_inject
#define GS_UTOFU_PUT_FLAGS_BASE
static utofu_vcq_hdl_t * gs_utofu_vcq
static size_t gs_utofu_ntni
static size_t gs_utofu_nalltni
static int gs_utofu_recv_rr
static unsigned long gs_utofu_put_flags
#define GS_UTOFU_PEND_STRIDE
void gs_utofu_post_puts(void *vctx, int tid, int nteam, int nsend, const int64_t *send_off, const int64_t *send_len, const uint64_t *rmt_vcq_id, const uint64_t *rmt_stadd, const int64_t *dst_off0, const int64_t *dst_off1, const int *dst_tag, int parity, int64_t send_base, int elem_size, int *ierr)
void gs_utofu_ctx_free(void *vctx, int *ierr)
static void gs_utofu_drain_tcq(size_t k)
void gs_utofu_poll_recv(void *vctx, int parity, int cap, int *ncompleted, int *out_idx, int *ierr)
void gs_utofu_ctx_create(void *send_buf, size_t send_bytes, void *recv_buf, size_t recv_bytes, void **ctx_out, int *nrvcq_out, uint64_t *recv_vcq_id, uint64_t *recv_stadd, int *ierr)
static long * gs_utofu_pend_slot(gs_utofu_ctx_t *ctx, size_t k, int half)
static size_t gs_utofu_nrvcq_def
static utofu_tni_id_t * gs_utofu_alltni
static size_t gs_utofu_nvcq
void gs_utofu_drain_half(void *vctx, int half, int tid, int nteam, int *ierr)
void gs_utofu_init(int ntni_req, int nthreads, int *ntni_out, int *nvcq_out, int *nrvcq_out, uint64_t *edata_max, int *ierr)
static uint64_t gs_utofu_edata_max
static int gs_utofu_put_one(gs_utofu_ctx_t *ctx, size_t k, int i, const int64_t *send_off, const int64_t *send_len, const uint64_t *rmt_vcq_id, const uint64_t *rmt_stadd, const int64_t *dst_off0, const int64_t *dst_off1, const int *dst_tag, int parity, int64_t send_base, int elem_size)
utofu_stadd_t * recv_stadd
utofu_vcq_hdl_t * recv_vcq
utofu_stadd_t * send_stadd
uint64_t * stash