88#define GS_UTOFU_PUT_FLAGS_BASE \
89 (UTOFU_ONESIDED_FLAG_TCQ_NOTICE | UTOFU_ONESIDED_FLAG_REMOTE_MRQ_NOTICE)
107#define GS_UTOFU_PEND_STRIDE 32
214 const char *
nv =
getenv(
"NEKO_GS_UTOFU_NVCQ");
255 const char *
nr =
getenv(
"NEKO_GS_UTOFU_NRVCQ");
266 size_t eb =
caps->max_edata_size;
274 const char *
ci =
getenv(
"NEKO_GS_UTOFU_CACHE_INJECT");
282 const char *mi =
getenv(
"NEKO_GS_UTOFU_MASTER_INJECT");
348 if (ctx->
nrvcq == 0) {
360 for (
size_t r = 0; r < ctx->
nrvcq; r++) {
368 for (
size_t r = 0; r < ctx->
nrvcq; r++) {
386 if (ctx ==
NULL)
return;
391 for (
int h = 0; h < 2; h++)
398 for (
size_t r = 0; r < ctx->
nrvcq; r++) {
430 size_t len = (
size_t)(send_len[
i] * elem_size);
468 int elem_size,
int *
ierr)
474 if (
tid != 0)
return;
478 rmt_stadd, dst_off0, dst_off1, dst_tag, parity,
492 rmt_vcq_id, rmt_stadd, dst_off0, dst_off1, dst_tag,
518 for (
int s = 0; s < ctx->
stash_n; s++) {
520 if ((
int)(
ed & 1u) == parity && n <
cap)
532 for (
size_t r = 0; r < ctx->
nrvcq && n <
cap; r++) {
544 if ((
int)(
ed & 1u) == parity)
__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
__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