From 4200c77d10cf0122d98cc6a357036a8c90b9ca86 Mon Sep 17 00:00:00 2001 From: Maxwell-Rosen Date: Wed, 26 Aug 2026 19:52:44 -0400 Subject: [PATCH 1/7] Add absolute value reduction functions for gkyl_array - Implement gkyl_array_reduce_range_sum_abs for CPU. - Implement gkyl_array_reduce_range_sum_abs_cu for CUDA. - Update tests to include absolute value reduction cases. - Modify gk_species to use absolute value reduction for diagnostics. --- core/unit/ctest_array_reduce.c | 52 ++++++++++++++++++++++++ core/zero/array_reduce.c | 25 ++++++++++++ core/zero/array_reduce_cu.cu | 46 +++++++++++++++++++++ core/zero/gkyl_array_reduce.h | 9 ++++ core/zero/gkyl_array_reduce_priv.h | 10 +++++ gyrokinetic/apps/gk_species.c | 8 ++-- gyrokinetic/apps/gkyl_gyrokinetic_priv.h | 2 +- 7 files changed, 148 insertions(+), 4 deletions(-) diff --git a/core/unit/ctest_array_reduce.c b/core/unit/ctest_array_reduce.c index 073d6b1508..60bd98f341 100644 --- a/core/unit/ctest_array_reduce.c +++ b/core/unit/ctest_array_reduce.c @@ -102,9 +102,59 @@ void test_sum_reduce_range() gkyl_array_release(arr); } +void test_sum_abs_reduce_range() +{ + int shape[] = {2, 3}; + struct gkyl_range range; + gkyl_range_init_from_shape(&range, 2, shape); + + struct gkyl_array *arr = gkyl_array_new(GKYL_DOUBLE, 2, range.volume); + for (size_t i=0; isize; ++i) { + double *d = gkyl_array_fetch(arr, i); + d[0] = i%2 == 0 ? 0.5 : -0.5; + d[1] = i%2 == 0 ? -1.5 : 1.5; + } + + double asum[2]; + gkyl_array_reduce_range_sum_abs(asum, arr, &range); + + TEST_CHECK( asum[0] == 0.5*range.volume ); + TEST_CHECK( asum[1] == 1.5*range.volume ); + + gkyl_array_release(arr); +} + // CUDA specific tests #ifdef GKYL_HAVE_CUDA +void test_cu_sum_abs_reduce_range() +{ + int shape[] = {2, 3}; + struct gkyl_range range; + gkyl_range_init_from_shape(&range, 2, shape); + + struct gkyl_array *arr = gkyl_array_new(GKYL_DOUBLE, 2, range.volume); + struct gkyl_array *arr_cu = gkyl_array_cu_dev_new(GKYL_DOUBLE, 2, range.volume); + for (size_t i=0; isize; ++i) { + double *d = gkyl_array_fetch(arr, i); + d[0] = i%2 == 0 ? 0.5 : -0.5; + d[1] = i%2 == 0 ? -1.5 : 1.5; + } + gkyl_array_copy(arr_cu, arr); + + double asum[2]; + double *asum_cu = gkyl_cu_malloc(2*sizeof(double)); + gkyl_array_reduce_range_sum_abs(asum_cu, arr_cu, &range); + gkyl_cu_memcpy(asum, asum_cu, sizeof(asum), GKYL_CU_MEMCPY_D2H); + + TEST_CHECK( asum[0] == 0.5*range.volume ); + TEST_CHECK( asum[1] == 1.5*range.volume ); + + gkyl_cu_free(asum_cu); + gkyl_array_release(arr_cu); + gkyl_array_release(arr); +} + void test_cu_array_reduce_max() { @@ -374,7 +424,9 @@ TEST_LIST = { { "array_reduce", test_reduce }, { "array_reduce_range", test_reduce_range }, { "array_reduce_sum_range", test_sum_reduce_range }, + { "array_reduce_sum_abs_range", test_sum_abs_reduce_range }, #ifdef GKYL_HAVE_CUDA + { "cu_array_reduce_sum_abs_range", test_cu_sum_abs_reduce_range }, { "cu_array_reduce_max", test_cu_array_reduce_max }, { "cu_array_reduce_max_big", test_cu_array_reduce_max_big }, { "cu_array_reduce_range_1d_max", test_cu_array_reduce_range_1d_max }, diff --git a/core/zero/array_reduce.c b/core/zero/array_reduce.c index 35f7740b99..5890711f53 100644 --- a/core/zero/array_reduce.c +++ b/core/zero/array_reduce.c @@ -119,3 +119,28 @@ gkyl_array_reduce_range(double *res, } } +void +gkyl_array_reduce_range_sum_abs(double *res, + const struct gkyl_array *arr, const struct gkyl_range *range) +{ + assert(arr->type == GKYL_DOUBLE); + +#ifdef GKYL_HAVE_CUDA + if (gkyl_array_is_cu_dev(arr)) { + gkyl_array_reduce_range_sum_abs_cu(res, arr, range); + return; + } +#endif + + long n = arr->ncomp; + for (long i=0; i +__global__ void +arraySumAbs_range_blockRedAtomic_cub(const struct gkyl_array* inp, + const struct gkyl_range range, double* out) +{ + unsigned long linc = blockIdx.x*blockDim.x + threadIdx.x; + + typedef cub::BlockReduce BlockReduceT; + __shared__ typename BlockReduceT::TempStorage temp; + + long nCells = range.volume; + size_t nComp = inp->ncomp; + int idx[GKYL_MAX_DIM]; + + for (size_t k = 0; k < nComp; ++k) { + double f = 0; + if (linc < nCells) { + gkyl_sub_range_inv_idx(&range, linc, idx); + long start = gkyl_range_idx(&range, idx); + const double* fptr = (const double*) gkyl_array_cfetch(inp, start); + f = fabs(fptr[k]); + } + double bResult = BlockReduceT(temp).Reduce(f, +#if CUDART_VERSION > 12090 + ::cuda::std::plus() +#else + cub::Sum() +#endif + ); + if (threadIdx.x == 0) + atomicAdd(&out[k], bResult); + } +} + +void +gkyl_array_reduce_range_sum_abs_cu(double *out_d, const struct gkyl_array* inp, + const struct gkyl_range *range) +{ + gkyl_cu_memset(out_d, 0, inp->ncomp*sizeof(double)); + + const int nthreads = GKYL_DEFAULT_NUM_THREADS; + int nblocks = gkyl_int_div_up(range->volume, nthreads); + arraySumAbs_range_blockRedAtomic_cub<<>>(inp->on_dev, *range, out_d); + // Device synchronize required because out_d may be host pinned memory. + cudaDeviceSynchronize(); +} diff --git a/core/zero/gkyl_array_reduce.h b/core/zero/gkyl_array_reduce.h index 93d97211a6..0f6fba5acd 100644 --- a/core/zero/gkyl_array_reduce.h +++ b/core/zero/gkyl_array_reduce.h @@ -25,3 +25,12 @@ void gkyl_array_reduce(double *res, const struct gkyl_array *arr, enum gkyl_arra void gkyl_array_reduce_range(double *res, const struct gkyl_array *arr, enum gkyl_array_op op, const struct gkyl_range *range); +/** + * Sum the absolute values of array data component-wise over a range. + * + * @param res On output, sums of absolute values (ncomp size). + * @param arr Array to perform reduction on. + * @param range Range specifying region. + */ +void gkyl_array_reduce_range_sum_abs(double *res, + const struct gkyl_array *arr, const struct gkyl_range *range); diff --git a/core/zero/gkyl_array_reduce_priv.h b/core/zero/gkyl_array_reduce_priv.h index cf032e450a..997739daec 100644 --- a/core/zero/gkyl_array_reduce_priv.h +++ b/core/zero/gkyl_array_reduce_priv.h @@ -55,4 +55,14 @@ void gkyl_array_reduce_sum_cu(double *out_d, const struct gkyl_array* inp); */ void gkyl_array_reduce_range_sum_cu(double *out_d, const struct gkyl_array* inp, const struct gkyl_range *range); +/** + * Sum the absolute values of a gkyl_array component-wise over a specified range. + * + * @param out_d A device array with as many elements as the 'inp' has components. + * @param inp A gkyl_array to be reduced. + * @param range A gkyl_range over which to perform the reduction. + */ +void gkyl_array_reduce_range_sum_abs_cu(double *out_d, const struct gkyl_array* inp, + const struct gkyl_range *range); + #endif diff --git a/gyrokinetic/apps/gk_species.c b/gyrokinetic/apps/gk_species.c index 2b93be63cf..e2706e7675 100644 --- a/gyrokinetic/apps/gk_species.c +++ b/gyrokinetic/apps/gk_species.c @@ -495,9 +495,11 @@ gk_species_calc_integrated_mom_dynamic(gkyl_gyrokinetic_app* app, struct gk_spec gkyl_dynvec_append(gks->integ_diag, tm, avals_global); if (gks->info.time_rate_diagnostics) { - // Reduce (sum) over whole domain, append to diagnostics. + // Sum the absolute time-rate moment in each cell over the whole domain, + // then append it to diagnostics. Taking the absolute value before the + // reduction prevents spatial cancellation. gkyl_array_accumulate(gks->fdot_mom_new, -1.0, gks->fdot_mom_old); - gkyl_array_reduce_range(gks->red_integ_diag, gks->fdot_mom_new, GKYL_SUM, &app->local); + gkyl_array_reduce_range_sum_abs(gks->red_integ_diag, gks->fdot_mom_new, &app->local); gkyl_comm_allreduce(app->comm, GKYL_DOUBLE, GKYL_SUM, num_mom, gks->red_integ_diag, gks->red_integ_diag_global); if (app->use_gpu) { @@ -562,7 +564,7 @@ gk_species_write_integrated_mom_dynamic(gkyl_gyrokinetic_app *app, struct gk_spe if (gks->is_first_fdot_integ_write_call) { struct gkyl_msgpack_map_elem io_meta_phi[] = { - { .key = "Description", .elem_type = GKYL_MP_STRING, .cval = "Volume integrated moments of time rate of change." } + { .key = "Description", .elem_type = GKYL_MP_STRING, .cval = "Volume integral of absolute moments of time rate of change." } }; int io_meta_len[] = {gks->io_meta_basic_len, app->gk_geom->io_meta_basic_len, 1}; const struct gkyl_msgpack_map_elem* io_meta[] = {gks->io_meta_basic, app->gk_geom->io_meta_basic, io_meta_phi}; diff --git a/gyrokinetic/apps/gkyl_gyrokinetic_priv.h b/gyrokinetic/apps/gkyl_gyrokinetic_priv.h index 5615160973..c88a47cbdb 100644 --- a/gyrokinetic/apps/gkyl_gyrokinetic_priv.h +++ b/gyrokinetic/apps/gkyl_gyrokinetic_priv.h @@ -1048,7 +1048,7 @@ struct gk_species { bool is_first_integ_write_call; // Whether dynvec is being written for the first time. struct gkyl_array *fdot_mom_old, *fdot_mom_new; // Moments of f_old and f_new. - gkyl_dynvec fdot_integ_diag; // Integrated moments of Delta f=f_new - f_old.. + gkyl_dynvec fdot_integ_diag; // Integrals of absolute moments of (f_new-f_old)/dt. bool is_first_fdot_integ_write_call; // Whether dynvec is being written for the first time. struct gkyl_array_integrate* integ_wfsq_op; // Operator to integrate w*f^2. From 1dc9b225dcae600f4b61b6b500caa4d0d2e067ef Mon Sep 17 00:00:00 2001 From: Maxwell-Rosen Date: Thu, 27 Aug 2026 10:01:47 -0400 Subject: [PATCH 2/7] Enhance diagnostics by adding absolute moment integration and updating related structures --- gyrokinetic/apps/gk_species.c | 49 +++++++++++++++++++++--- gyrokinetic/apps/gkyl_gyrokinetic_priv.h | 4 +- gyrokinetic/apps/gyrokinetic.c | 4 +- 3 files changed, 49 insertions(+), 8 deletions(-) diff --git a/gyrokinetic/apps/gk_species.c b/gyrokinetic/apps/gk_species.c index e2706e7675..0b0c88babf 100644 --- a/gyrokinetic/apps/gk_species.c +++ b/gyrokinetic/apps/gk_species.c @@ -495,11 +495,10 @@ gk_species_calc_integrated_mom_dynamic(gkyl_gyrokinetic_app* app, struct gk_spec gkyl_dynvec_append(gks->integ_diag, tm, avals_global); if (gks->info.time_rate_diagnostics) { - // Sum the absolute time-rate moment in each cell over the whole domain, - // then append it to diagnostics. Taking the absolute value before the - // reduction prevents spatial cancellation. gkyl_array_accumulate(gks->fdot_mom_new, -1.0, gks->fdot_mom_old); - gkyl_array_reduce_range_sum_abs(gks->red_integ_diag, gks->fdot_mom_new, &app->local); + + // Sum the signed time-rate moment over the whole domain. + gkyl_array_reduce_range(gks->red_integ_diag, gks->fdot_mom_new, GKYL_SUM, &app->local); gkyl_comm_allreduce(app->comm, GKYL_DOUBLE, GKYL_SUM, num_mom, gks->red_integ_diag, gks->red_integ_diag_global); if (app->use_gpu) { @@ -509,6 +508,19 @@ gk_species_calc_integrated_mom_dynamic(gkyl_gyrokinetic_app* app, struct gk_spec memcpy(avals_global, gks->red_integ_diag_global, sizeof(double[num_mom])); } gkyl_dynvec_append(gks->fdot_integ_diag, tm, avals_global); + + // Sum the absolute time-rate moment in each cell over the whole domain. + // Taking the absolute value before the reduction prevents spatial cancellation. + gkyl_array_reduce_range_sum_abs(gks->red_integ_diag, gks->fdot_mom_new, &app->local); + gkyl_comm_allreduce(app->comm, GKYL_DOUBLE, GKYL_SUM, num_mom, + gks->red_integ_diag, gks->red_integ_diag_global); + if (app->use_gpu) { + gkyl_cu_memcpy(avals_global, gks->red_integ_diag_global, sizeof(double[num_mom]), GKYL_CU_MEMCPY_D2H); + } + else { + memcpy(avals_global, gks->red_integ_diag_global, sizeof(double[num_mom])); + } + gkyl_dynvec_append(gks->fdot_abs_integ_diag, tm, avals_global); } app->stat.species_diag_calc_tm += gkyl_time_diff_now_sec(wst); @@ -564,7 +576,7 @@ gk_species_write_integrated_mom_dynamic(gkyl_gyrokinetic_app *app, struct gk_spe if (gks->is_first_fdot_integ_write_call) { struct gkyl_msgpack_map_elem io_meta_phi[] = { - { .key = "Description", .elem_type = GKYL_MP_STRING, .cval = "Volume integral of absolute moments of time rate of change." } + { .key = "Description", .elem_type = GKYL_MP_STRING, .cval = "Volume integrated moments of time rate of change." } }; int io_meta_len[] = {gks->io_meta_basic_len, app->gk_geom->io_meta_basic_len, 1}; const struct gkyl_msgpack_map_elem* io_meta[] = {gks->io_meta_basic, app->gk_geom->io_meta_basic, io_meta_phi}; @@ -577,9 +589,31 @@ gk_species_write_integrated_mom_dynamic(gkyl_gyrokinetic_app *app, struct gk_spe else { gkyl_dynvec_awrite(gks->fdot_integ_diag, fileNm); } + + const char *abs_fmt = "%s-%s_fdot_abs_%s.gkyl"; + int abs_sz = gkyl_calc_strlen(abs_fmt, app->name, gks->info.name, "integrated_moms"); + char abs_fileNm[abs_sz+1]; // ensures no buffer overflow + snprintf(abs_fileNm, sizeof abs_fileNm, abs_fmt, app->name, gks->info.name, "integrated_moms"); + + if (gks->is_first_fdot_abs_integ_write_call) { + struct gkyl_msgpack_map_elem io_meta_phi[] = { + { .key = "Description", .elem_type = GKYL_MP_STRING, .cval = "Volume integral of absolute moments of time rate of change." } + }; + int io_meta_len[] = {gks->io_meta_basic_len, app->gk_geom->io_meta_basic_len, 1}; + const struct gkyl_msgpack_map_elem* io_meta[] = {gks->io_meta_basic, app->gk_geom->io_meta_basic, io_meta_phi}; + struct gkyl_msgpack_data *mt = gkyl_msgpack_create_union(sizeof(io_meta_len)/sizeof(int), io_meta_len, io_meta); + + gkyl_dynvec_write_wmeta(gks->fdot_abs_integ_diag, abs_fileNm, mt); + gks->is_first_fdot_abs_integ_write_call = false; + gkyl_msgpack_data_release(mt); + } + else { + gkyl_dynvec_awrite(gks->fdot_abs_integ_diag, abs_fileNm); + } } gkyl_dynvec_clear(gks->fdot_integ_diag); - app->stat.n_diag_io += 1; + gkyl_dynvec_clear(gks->fdot_abs_integ_diag); + app->stat.n_diag_io += 2; } app->stat.species_diag_io_tm += gkyl_time_diff_now_sec(wst); @@ -750,6 +784,7 @@ gk_species_release_dynamic(const gkyl_gyrokinetic_app* app, const struct gk_spec gkyl_array_release(s->fdot_mom_old); gkyl_array_release(s->fdot_mom_new); gkyl_dynvec_release(s->fdot_integ_diag); + gkyl_dynvec_release(s->fdot_abs_integ_diag); } } @@ -812,7 +847,9 @@ gk_species_init_dynamic(struct gkyl_gk *gk_app_inp, struct gkyl_gyrokinetic_app gks->fdot_mom_old = mkarr(app->use_gpu, gks->integ_moms.marr->ncomp, gks->integ_moms.marr->size); gks->fdot_mom_new = mkarr(app->use_gpu, gks->integ_moms.marr->ncomp, gks->integ_moms.marr->size); gks->fdot_integ_diag = gkyl_dynvec_new(GKYL_DOUBLE, gks->integ_moms.num_mom); + gks->fdot_abs_integ_diag = gkyl_dynvec_new(GKYL_DOUBLE, gks->integ_moms.num_mom); gks->is_first_fdot_integ_write_call = true; + gks->is_first_fdot_abs_integ_write_call = true; } // Objects for L2 norm diagnostic. diff --git a/gyrokinetic/apps/gkyl_gyrokinetic_priv.h b/gyrokinetic/apps/gkyl_gyrokinetic_priv.h index c88a47cbdb..3e8f86c231 100644 --- a/gyrokinetic/apps/gkyl_gyrokinetic_priv.h +++ b/gyrokinetic/apps/gkyl_gyrokinetic_priv.h @@ -1048,8 +1048,10 @@ struct gk_species { bool is_first_integ_write_call; // Whether dynvec is being written for the first time. struct gkyl_array *fdot_mom_old, *fdot_mom_new; // Moments of f_old and f_new. - gkyl_dynvec fdot_integ_diag; // Integrals of absolute moments of (f_new-f_old)/dt. + gkyl_dynvec fdot_integ_diag; // Integrated moments of (f_new-f_old)/dt. + gkyl_dynvec fdot_abs_integ_diag; // Integrals of absolute moments of (f_new-f_old)/dt. bool is_first_fdot_integ_write_call; // Whether dynvec is being written for the first time. + bool is_first_fdot_abs_integ_write_call; // Whether absolute fdot dynvec is being written for the first time. struct gkyl_array_integrate* integ_wfsq_op; // Operator to integrate w*f^2. double *L2norm_local, *L2norm_global; // L2norm in local MPI process and across the communicator. diff --git a/gyrokinetic/apps/gyrokinetic.c b/gyrokinetic/apps/gyrokinetic.c index 526aeebc67..baf7adf3c8 100644 --- a/gyrokinetic/apps/gyrokinetic.c +++ b/gyrokinetic/apps/gyrokinetic.c @@ -3390,8 +3390,10 @@ gkyl_gyrokinetic_app_from_frame_species(gkyl_gyrokinetic_app *app, int sidx, int gk_s->is_first_L2norm_write_call = false; for (int b=0; bbflux.num_boundaries; ++b) gk_s->bflux.is_first_intmom_write_call[b] = false; - if (gk_s->info.time_rate_diagnostics) + if (gk_s->info.time_rate_diagnostics) { gk_s->is_first_fdot_integ_write_call = false; + gk_s->is_first_fdot_abs_integ_write_call = false; + } if (gk_s->positivity.type) gk_s->positivity.is_first_integ_write_call = false; if (gk_s->rad.radiation_id == GKYL_GK_RADIATION) From f48e4601bb012dc8a7f6da681b2656a2158a2cea Mon Sep 17 00:00:00 2001 From: Maxwell-Rosen Date: Mon, 31 Aug 2026 13:28:21 -0400 Subject: [PATCH 3/7] Enhance time rate diagnostics across multiple gyrokinetic modules - Updated the time rate diagnostics configuration in various files to support two diagnostic types: - GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS - GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS - Replaced the previous boolean flag for time rate diagnostics with a structured approach that includes the number of diagnostics and their specific types. - This change improves the diagnostic capabilities of the gyrokinetic simulations, allowing for more detailed analysis of the integrated moments over time. --- gyrokinetic/apps/gk_species.c | 167 ++++++++++++++---- gyrokinetic/apps/gkyl_gyrokinetic.h | 12 +- gyrokinetic/apps/gkyl_gyrokinetic_lw.h | 8 + gyrokinetic/apps/gkyl_gyrokinetic_multib.h | 4 +- gyrokinetic/apps/gkyl_gyrokinetic_priv.h | 16 ++ gyrokinetic/apps/gyrokinetic.c | 4 +- gyrokinetic/apps/gyrokinetic_lw.c | 25 +++ gyrokinetic/apps/gyrokinetic_multib.c | 5 +- .../apps/gyrokinetic_multib_update_ssp_rk3.c | 3 +- gyrokinetic/apps/gyrokinetic_update_ssp_rk3.c | 2 + .../creg/rt_gk_asdex_solonly_3x2v_p1.c | 12 +- gyrokinetic/creg/rt_gk_cbc_2x2v_p1.c | 12 +- gyrokinetic/creg/rt_gk_cbc_3x2v_p1.c | 12 +- gyrokinetic/creg/rt_gk_d3d_iwl_2x2v_p1.c | 12 +- gyrokinetic/creg/rt_gk_d3d_iwl_3x2v_p1.c | 12 +- gyrokinetic/creg/rt_gk_helical_zpar_3x2v_p1.c | 12 +- .../creg/rt_gk_helical_zvert_3x2v_p1.c | 12 +- gyrokinetic/creg/rt_gk_lapd_cyl_3x2v_p1.c | 12 +- gyrokinetic/creg/rt_gk_leaky_bag_1x2v_p1.c | 6 +- .../creg/rt_gk_leaky_bag_1x2v_p1_nux_center.c | 6 +- .../creg/rt_gk_leaky_bag_1x2v_p1_nux_edge.c | 6 +- .../creg/rt_gk_leaky_bag_1x2v_p1_ux_stretch.c | 6 +- .../rt_gk_leaky_bag_boltz_sheath_1x2v_p1.c | 7 +- ...eaky_bag_boltz_sheath_1x2v_p1_nux_center.c | 6 +- ..._leaky_bag_boltz_sheath_1x2v_p1_nux_edge.c | 6 +- ...eaky_bag_boltz_sheath_1x2v_p1_ux_stretch.c | 6 +- gyrokinetic/creg/rt_gk_multib_asdex_2x2v_p1.c | 12 +- .../creg/rt_gk_multib_asdex_solonly_3x2v_p1.c | 12 +- .../creg/rt_gk_multib_nstx_solonly_3x2v_p1.c | 12 +- .../creg/rt_gk_multib_sheath_1x2v_p1.c | 12 +- gyrokinetic/creg/rt_gk_multib_slab_2x2v_p1.c | 12 +- gyrokinetic/creg/rt_gk_multib_step_2x2v_p1.c | 12 +- .../creg/rt_gk_multib_step_eirene_2x2v_p1.c | 12 +- .../rt_gk_multib_step_nonuniform_2x2v_p1.c | 12 +- .../creg/rt_gk_multib_step_sol_2x2v_p1.c | 12 +- .../creg/rt_gk_multib_tcv_x21_3x2v_p1.c | 12 +- gyrokinetic/creg/rt_gk_neut_sheath_3x2v_p1.c | 12 +- gyrokinetic/creg/rt_gk_sheath_1x2v_p1.c | 12 +- .../creg/rt_gk_sheath_1x2v_p1_ic_import.c | 12 +- gyrokinetic/creg/rt_gk_sheath_2x2v_p1.c | 12 +- gyrokinetic/creg/rt_gk_sheath_3x2v_p1.c | 12 +- .../creg/rt_gk_sheath_fluid_neut_1x2v_p1.c | 12 +- .../creg/rt_gk_sheath_heat_source_2x2v_p1.c | 12 +- gyrokinetic/creg/rt_gk_solovev_out_3x2v_p1.c | 12 +- gyrokinetic/creg/rt_gk_step_2x2v_p1_cons.c | 12 +- gyrokinetic/creg/rt_gk_step_out_2x2v_p1.c | 12 +- gyrokinetic/creg/rt_gk_tcv_core_3x2v_p1.c | 12 +- .../creg/rt_gk_tcv_iwl_adapt_source_2x2v_p1.c | 12 +- .../creg/rt_gk_tcv_iwl_adapt_source_3x2v_p1.c | 12 +- gyrokinetic/creg/rt_gk_wham_1x2v_p1.c | 12 +- .../creg/rt_gk_wham_boltz_elc_poa_1x2v_p1.c | 6 +- 51 files changed, 576 insertions(+), 109 deletions(-) diff --git a/gyrokinetic/apps/gk_species.c b/gyrokinetic/apps/gk_species.c index 0b0c88babf..a0d7490376 100644 --- a/gyrokinetic/apps/gk_species.c +++ b/gyrokinetic/apps/gk_species.c @@ -348,6 +348,36 @@ gk_species_write_dynamic(gkyl_gyrokinetic_app* app, struct gk_species *gks, doub app->stat.species_io_tm += gkyl_time_diff_now_sec(wst); app->stat.n_io += 1; + + if (gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT]) { + struct timespec wdt = gkyl_wall_clock(); + + struct gkyl_msgpack_map_elem io_meta_fdot[] = { + { .key = "Description", .elem_type = GKYL_MP_STRING, + .cval = "Time rate of change of the distribution function." } + }; + int io_meta_fdot_len[] = {gks->io_meta_phase_len, app->gk_geom->io_meta_basic_len, 1}; + const struct gkyl_msgpack_map_elem* io_meta_fdot_union[] = { + gks->io_meta_phase, app->gk_geom->io_meta_basic, io_meta_fdot + }; + struct gkyl_msgpack_data *fdot_mt = gkyl_msgpack_create_union( + sizeof(io_meta_fdot_len)/sizeof(int), io_meta_fdot_len, io_meta_fdot_union); + + const char *fdot_fmt = "%s-%s_fdot_%d.gkyl"; + int fdot_sz = gkyl_calc_strlen(fdot_fmt, app->name, gks->info.name, frame); + char fdot_fileNm[fdot_sz+1]; // Ensures no buffer overflow. + snprintf(fdot_fileNm, sizeof fdot_fileNm, fdot_fmt, app->name, gks->info.name, frame); + + if (app->use_gpu) { + gkyl_array_copy(gks->fdot_host, gks->fdot); + } + gkyl_comm_array_write(gks->comm, &gks->grid, &gks->local, fdot_mt, + gks->fdot_host, fdot_fileNm); + + gkyl_msgpack_data_release(fdot_mt); + app->stat.species_diag_io_tm += gkyl_time_diff_now_sec(wdt); + app->stat.n_diag_io += 1; + } } static void @@ -471,6 +501,26 @@ gk_species_calc_int_mom_dt(gkyl_gyrokinetic_app* app, struct gk_species *gks, do gks->calc_int_mom_dt_func(app, gks, dt, fdot_int_mom); } +void +gk_species_calc_fdot_begin_step(gkyl_gyrokinetic_app* app, struct gk_species *gks, double dt) +{ + if (gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT]) { + struct timespec wst = gkyl_wall_clock(); + gkyl_array_set(gks->fdot, -1.0/dt, gks->f); + app->stat.fdot_tm += gkyl_time_diff_now_sec(wst); + } +} + +void +gk_species_calc_fdot_complete_step(gkyl_gyrokinetic_app* app, struct gk_species *gks, double dt) +{ + if (gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT]) { + struct timespec wst = gkyl_wall_clock(); + gkyl_array_accumulate(gks->fdot, 1.0/dt, gks->f); + app->stat.fdot_tm += gkyl_time_diff_now_sec(wst); + } +} + static void gk_species_calc_integrated_mom_dynamic(gkyl_gyrokinetic_app* app, struct gk_species *gks, double tm) { @@ -494,33 +544,41 @@ gk_species_calc_integrated_mom_dynamic(gkyl_gyrokinetic_app* app, struct gk_spec } gkyl_dynvec_append(gks->integ_diag, tm, avals_global); - if (gks->info.time_rate_diagnostics) { + bool calc_fdot_integ = + gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS]; + bool calc_fdot_abs_integ = + gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS]; + if (calc_fdot_integ || calc_fdot_abs_integ) { gkyl_array_accumulate(gks->fdot_mom_new, -1.0, gks->fdot_mom_old); - // Sum the signed time-rate moment over the whole domain. - gkyl_array_reduce_range(gks->red_integ_diag, gks->fdot_mom_new, GKYL_SUM, &app->local); - gkyl_comm_allreduce(app->comm, GKYL_DOUBLE, GKYL_SUM, num_mom, - gks->red_integ_diag, gks->red_integ_diag_global); - if (app->use_gpu) { - gkyl_cu_memcpy(avals_global, gks->red_integ_diag_global, sizeof(double[num_mom]), GKYL_CU_MEMCPY_D2H); - } - else { - memcpy(avals_global, gks->red_integ_diag_global, sizeof(double[num_mom])); + if (calc_fdot_integ) { + // Sum the signed time-rate moment over the whole domain. + gkyl_array_reduce_range(gks->red_integ_diag, gks->fdot_mom_new, GKYL_SUM, &app->local); + gkyl_comm_allreduce(app->comm, GKYL_DOUBLE, GKYL_SUM, num_mom, + gks->red_integ_diag, gks->red_integ_diag_global); + if (app->use_gpu) { + gkyl_cu_memcpy(avals_global, gks->red_integ_diag_global, sizeof(double[num_mom]), GKYL_CU_MEMCPY_D2H); + } + else { + memcpy(avals_global, gks->red_integ_diag_global, sizeof(double[num_mom])); + } + gkyl_dynvec_append(gks->fdot_integ_diag, tm, avals_global); } - gkyl_dynvec_append(gks->fdot_integ_diag, tm, avals_global); - // Sum the absolute time-rate moment in each cell over the whole domain. - // Taking the absolute value before the reduction prevents spatial cancellation. - gkyl_array_reduce_range_sum_abs(gks->red_integ_diag, gks->fdot_mom_new, &app->local); - gkyl_comm_allreduce(app->comm, GKYL_DOUBLE, GKYL_SUM, num_mom, - gks->red_integ_diag, gks->red_integ_diag_global); - if (app->use_gpu) { - gkyl_cu_memcpy(avals_global, gks->red_integ_diag_global, sizeof(double[num_mom]), GKYL_CU_MEMCPY_D2H); - } - else { - memcpy(avals_global, gks->red_integ_diag_global, sizeof(double[num_mom])); + if (calc_fdot_abs_integ) { + // Sum the absolute time-rate moment in each cell over the whole domain. + // Taking the absolute value before the reduction prevents spatial cancellation. + gkyl_array_reduce_range_sum_abs(gks->red_integ_diag, gks->fdot_mom_new, &app->local); + gkyl_comm_allreduce(app->comm, GKYL_DOUBLE, GKYL_SUM, num_mom, + gks->red_integ_diag, gks->red_integ_diag_global); + if (app->use_gpu) { + gkyl_cu_memcpy(avals_global, gks->red_integ_diag_global, sizeof(double[num_mom]), GKYL_CU_MEMCPY_D2H); + } + else { + memcpy(avals_global, gks->red_integ_diag_global, sizeof(double[num_mom])); + } + gkyl_dynvec_append(gks->fdot_abs_integ_diag, tm, avals_global); } - gkyl_dynvec_append(gks->fdot_abs_integ_diag, tm, avals_global); } app->stat.species_diag_calc_tm += gkyl_time_diff_now_sec(wst); @@ -566,7 +624,7 @@ gk_species_write_integrated_mom_dynamic(gkyl_gyrokinetic_app *app, struct gk_spe gkyl_dynvec_clear(gks->integ_diag); app->stat.n_diag_io += 1; - if (gks->info.time_rate_diagnostics) { + if (gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS]) { if (rank == 0) { // Write integrated diagnostic moments. const char *fmt = "%s-%s_fdot_%s.gkyl"; @@ -590,6 +648,13 @@ gk_species_write_integrated_mom_dynamic(gkyl_gyrokinetic_app *app, struct gk_spe gkyl_dynvec_awrite(gks->fdot_integ_diag, fileNm); } + } + gkyl_dynvec_clear(gks->fdot_integ_diag); + app->stat.n_diag_io += 1; + } + + if (gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS]) { + if (rank == 0) { const char *abs_fmt = "%s-%s_fdot_abs_%s.gkyl"; int abs_sz = gkyl_calc_strlen(abs_fmt, app->name, gks->info.name, "integrated_moms"); char abs_fileNm[abs_sz+1]; // ensures no buffer overflow @@ -611,9 +676,8 @@ gk_species_write_integrated_mom_dynamic(gkyl_gyrokinetic_app *app, struct gk_spe gkyl_dynvec_awrite(gks->fdot_abs_integ_diag, abs_fileNm); } } - gkyl_dynvec_clear(gks->fdot_integ_diag); gkyl_dynvec_clear(gks->fdot_abs_integ_diag); - app->stat.n_diag_io += 2; + app->stat.n_diag_io += 1; } app->stat.species_diag_io_tm += gkyl_time_diff_now_sec(wst); @@ -779,11 +843,18 @@ gk_species_release_dynamic(const gkyl_gyrokinetic_app* app, const struct gk_spec gkyl_free(s->L2norm_global); } - if (s->info.time_rate_diagnostics) { - // Free df/dt diagnostics memory. + bool calc_fdot_integ = + s->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS]; + bool calc_fdot_abs_integ = + s->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS]; + if (calc_fdot_integ || calc_fdot_abs_integ) { gkyl_array_release(s->fdot_mom_old); gkyl_array_release(s->fdot_mom_new); + } + if (calc_fdot_integ) { gkyl_dynvec_release(s->fdot_integ_diag); + } + if (calc_fdot_abs_integ) { gkyl_dynvec_release(s->fdot_abs_integ_diag); } } @@ -842,13 +913,21 @@ gk_species_init_dynamic(struct gkyl_gk *gk_app_inp, struct gkyl_gyrokinetic_app gks->integ_diag = gkyl_dynvec_new(GKYL_DOUBLE, gks->integ_moms.num_mom); gks->is_first_integ_write_call = true; - // Allocate dynamic-vector to store Delta f integrated moments. - if (gks->info.time_rate_diagnostics) { + // Allocate dynamic-vectors to store Delta f integrated moments. + bool calc_fdot_integ = + gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS]; + bool calc_fdot_abs_integ = + gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS]; + if (calc_fdot_integ || calc_fdot_abs_integ) { gks->fdot_mom_old = mkarr(app->use_gpu, gks->integ_moms.marr->ncomp, gks->integ_moms.marr->size); gks->fdot_mom_new = mkarr(app->use_gpu, gks->integ_moms.marr->ncomp, gks->integ_moms.marr->size); + } + if (calc_fdot_integ) { gks->fdot_integ_diag = gkyl_dynvec_new(GKYL_DOUBLE, gks->integ_moms.num_mom); - gks->fdot_abs_integ_diag = gkyl_dynvec_new(GKYL_DOUBLE, gks->integ_moms.num_mom); gks->is_first_fdot_integ_write_call = true; + } + if (calc_fdot_abs_integ) { + gks->fdot_abs_integ_diag = gkyl_dynvec_new(GKYL_DOUBLE, gks->integ_moms.num_mom); gks->is_first_fdot_abs_integ_write_call = true; } @@ -1094,7 +1173,7 @@ gk_species_init_dynamic(struct gkyl_gk *gk_app_inp, struct gkyl_gyrokinetic_app gks->write_integrated_mom_func = gk_species_write_integrated_mom_dynamic; gks->calc_L2norm_func = gk_species_calc_L2norm_dynamic; gks->write_L2norm_func = gk_species_write_L2norm_dynamic; - if (gks->info.time_rate_diagnostics) + if (calc_fdot_integ || calc_fdot_abs_integ) gks->calc_int_mom_dt_func = gk_species_calc_int_mom_dt_enabled; else gks->calc_int_mom_dt_func = gk_species_calc_int_mom_dt_disabled; @@ -1421,6 +1500,17 @@ gk_species_init(struct gkyl_gk *gk_app_inp, struct gkyl_gyrokinetic_app *app, st int cdim = app->cdim, vdim = gks->info.vdim; int pdim = cdim+vdim; + assert(gks->info.num_time_rate_diagnostics >= 0); + assert(gks->info.num_time_rate_diagnostics <= GKYL_GK_TIME_RATE_DIAGNOSTIC_NUM); + for (int d=0; dtime_rate_diagnostics[d] = false; + } + for (int d=0; dinfo.num_time_rate_diagnostics; ++d) { + enum gkyl_gyrokinetic_time_rate_diagnostic diag = gks->info.time_rate_diagnostics[d]; + assert(diag >= 0 && diag < GKYL_GK_TIME_RATE_DIAGNOSTIC_NUM); + gks->time_rate_diagnostics[diag] = true; + } + int cells[GKYL_MAX_DIM], ghost[GKYL_MAX_DIM]; double lower[GKYL_MAX_DIM], upper[GKYL_MAX_DIM]; @@ -1583,6 +1673,14 @@ gk_species_init(struct gkyl_gk *gk_app_inp, struct gkyl_gyrokinetic_app *app, st if (app->use_gpu) gks->f_host = mkarr(false, gks->basis.num_basis, gks->local_ext.volume); + if (gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT]) { + gks->fdot = mkarr(app->use_gpu, gks->basis.num_basis, gks->local_ext.volume); + gks->fdot_host = gks->fdot; + if (app->use_gpu) { + gks->fdot_host = mkarr(false, gks->basis.num_basis, gks->local_ext.volume); + } + } + // Allocate cflrate (scalar array). gks->cflrate = mkarr(app->use_gpu, 1, gks->local_ext.volume); @@ -2054,6 +2152,13 @@ gk_species_release(const gkyl_gyrokinetic_app* app, const struct gk_species *gks gkyl_cu_free(gks->basis_on_dev); } + if (gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT]) { + gkyl_array_release(gks->fdot); + if (app->use_gpu) { + gkyl_array_release(gks->fdot_host); + } + } + gkyl_velocity_map_release(gks->vel_map); gk_species_collisionless_release(app, &gks->collisionless); diff --git a/gyrokinetic/apps/gkyl_gyrokinetic.h b/gyrokinetic/apps/gkyl_gyrokinetic.h index e04e1510ba..cf8ddff4cf 100644 --- a/gyrokinetic/apps/gkyl_gyrokinetic.h +++ b/gyrokinetic/apps/gkyl_gyrokinetic.h @@ -95,6 +95,14 @@ struct gkyl_phase_diagnostics_inp { bool time_integrated; // Whether to use time integrated diags. }; +// Time-rate diagnostics for a gyrokinetic distribution function. +enum gkyl_gyrokinetic_time_rate_diagnostic { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT = 0, // Phase-space (f_new-f_old)/dt. + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, // Volume-integrated moments of fdot. + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, // Volume integral of the absolute fdot moments. + GKYL_GK_TIME_RATE_DIAGNOSTIC_NUM, +}; + // Parameters for collisionless terms. struct gkyl_gyrokinetic_collisionless { enum gkyl_gk_collisionless_type type; // Type of collisionless terms. @@ -450,7 +458,9 @@ struct gkyl_gyrokinetic_species { enum gkyl_distribution_moments diag_moments[12]; // list of diagnostic moments int num_integrated_diag_moments; // Number of integrated diagnostic moments. enum gkyl_distribution_moments integrated_diag_moments[12]; // List of integrated diagnostic moments. - bool time_rate_diagnostics; // Whether to ouput df/dt diagnostics. + int num_time_rate_diagnostics; // Number of time-rate diagnostics. + enum gkyl_gyrokinetic_time_rate_diagnostic + time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_NUM]; // List of time-rate diagnostics. bool write_omega_cfl; // Whether to ouput dt diagnostic for the CFL constraint. struct gkyl_gyrokinetic_collisionless collisionless; // Collisionless terms. diff --git a/gyrokinetic/apps/gkyl_gyrokinetic_lw.h b/gyrokinetic/apps/gkyl_gyrokinetic_lw.h index 90bc66fcb8..d8d8ddbd04 100644 --- a/gyrokinetic/apps/gkyl_gyrokinetic_lw.h +++ b/gyrokinetic/apps/gkyl_gyrokinetic_lw.h @@ -45,6 +45,14 @@ gkyl_register_gyrokinetic_position_map_types(lua_State *L); void gkyl_register_gyrokinetic_field_types(lua_State *L); +/** + * Add time-rate diagnostic flags for gyrokinetic species into Lua interpreter. + * + * @param L Lua state to use. + */ +void +gkyl_register_gyrokinetic_time_rate_diagnostic_types(lua_State *L); + /** * Add radiation type flags for gyrokinetic species initialization into Lua interpreter. * diff --git a/gyrokinetic/apps/gkyl_gyrokinetic_multib.h b/gyrokinetic/apps/gkyl_gyrokinetic_multib.h index e47288d6f1..f3e6fbeeb3 100644 --- a/gyrokinetic/apps/gkyl_gyrokinetic_multib.h +++ b/gyrokinetic/apps/gkyl_gyrokinetic_multib.h @@ -32,7 +32,9 @@ struct gkyl_gyrokinetic_multib_species { enum gkyl_distribution_moments diag_moments[12]; // List of diagnostic moments. int num_integrated_diag_moments; // Number of integrated diagnostic moments. enum gkyl_distribution_moments integrated_diag_moments[12]; // List of integrated diagnostic moments. - bool time_rate_diagnostics; // Whether to ouput df/dt diagnostics. + int num_time_rate_diagnostics; // Number of time-rate diagnostics. + enum gkyl_gyrokinetic_time_rate_diagnostic + time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_NUM]; // List of time-rate diagnostics. struct gkyl_phase_diagnostics_inp boundary_flux_diagnostics; diff --git a/gyrokinetic/apps/gkyl_gyrokinetic_priv.h b/gyrokinetic/apps/gkyl_gyrokinetic_priv.h index 3e8f86c231..076736cb43 100644 --- a/gyrokinetic/apps/gkyl_gyrokinetic_priv.h +++ b/gyrokinetic/apps/gkyl_gyrokinetic_priv.h @@ -1030,6 +1030,7 @@ struct gk_species { int io_meta_conf_len; // Number of elements in io_meta_conf. struct gkyl_array *f, *f1, *fnew; // Arrays for updates. + struct gkyl_array *fdot, *fdot_host; // Phase-space (f_new-f_old)/dt and its host copy. struct gkyl_array *cflrate; // CFL rate in each cell. struct gkyl_array *cflrate_ho; // CFL rate in each cell on host-side. @@ -1053,6 +1054,9 @@ struct gk_species { bool is_first_fdot_integ_write_call; // Whether dynvec is being written for the first time. bool is_first_fdot_abs_integ_write_call; // Whether absolute fdot dynvec is being written for the first time. + // Lookup table populated from info.time_rate_diagnostics. + bool time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_NUM]; + struct gkyl_array_integrate* integ_wfsq_op; // Operator to integrate w*f^2. double *L2norm_local, *L2norm_global; // L2norm in local MPI process and across the communicator. gkyl_dynvec L2norm; // L2 norm. @@ -3334,6 +3338,18 @@ void gk_species_write_L2norm(gkyl_gyrokinetic_app* app, struct gk_species *gks); void gk_species_calc_int_mom_dt(gkyl_gyrokinetic_app* app, struct gk_species *gks, double dt, struct gkyl_array *fdot_int_mom); +/** + * Store the old-distribution contribution to the phase-space fdot diagnostic. + */ +void +gk_species_calc_fdot_begin_step(gkyl_gyrokinetic_app* app, struct gk_species *gks, double dt); + +/** + * Add the new-distribution contribution to the phase-space fdot diagnostic. + */ +void +gk_species_calc_fdot_complete_step(gkyl_gyrokinetic_app* app, struct gk_species *gks, double dt); + /** * Delete resources used in species. * diff --git a/gyrokinetic/apps/gyrokinetic.c b/gyrokinetic/apps/gyrokinetic.c index baf7adf3c8..bec23aeae7 100644 --- a/gyrokinetic/apps/gyrokinetic.c +++ b/gyrokinetic/apps/gyrokinetic.c @@ -3390,8 +3390,10 @@ gkyl_gyrokinetic_app_from_frame_species(gkyl_gyrokinetic_app *app, int sidx, int gk_s->is_first_L2norm_write_call = false; for (int b=0; bbflux.num_boundaries; ++b) gk_s->bflux.is_first_intmom_write_call[b] = false; - if (gk_s->info.time_rate_diagnostics) { + if (gk_s->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS]) { gk_s->is_first_fdot_integ_write_call = false; + } + if (gk_s->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS]) { gk_s->is_first_fdot_abs_integ_write_call = false; } if (gk_s->positivity.type) diff --git a/gyrokinetic/apps/gyrokinetic_lw.c b/gyrokinetic/apps/gyrokinetic_lw.c index 7a1a1500a7..4b3e2b7a4f 100644 --- a/gyrokinetic/apps/gyrokinetic_lw.c +++ b/gyrokinetic/apps/gyrokinetic_lw.c @@ -75,6 +75,14 @@ static const struct gkyl_str_int_pair gk_field_type[] = { { 0, 0 } }; +// Gyrokinetic distribution-function time-rate diagnostics -> enum map. +static const struct gkyl_str_int_pair gk_time_rate_diagnostic_type[] = { + { "Fdot", GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT }, + { "FdotIntegratedMoments", GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS }, + { "FdotAbsIntegratedMoments", GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS }, + { 0, 0 } +}; + // Gyrokinetic radiation type -> enum map. static const struct gkyl_str_int_pair gk_radiation_type[] = { { "None", GKYL_NO_RADIATION }, @@ -181,6 +189,12 @@ gkyl_register_gyrokinetic_field_types(lua_State *L) register_types(L, gk_field_type, "GKField"); } +void +gkyl_register_gyrokinetic_time_rate_diagnostic_types(lua_State *L) +{ + register_types(L, gk_time_rate_diagnostic_type, "TimeRateDiagnostic"); +} + void gkyl_register_gyrokinetic_radiation_types(lua_State *L) { @@ -389,6 +403,16 @@ gyrokinetic_species_lw_new(lua_State *L) gk_species.num_diag_moments = num_diag_moments; } + with_lua_tbl_tbl(L, "timeRateDiagnostics") { + int num_time_rate_diagnostics = glua_objlen(L); + + for (int i = 0; i < num_time_rate_diagnostics; i ++) { + gk_species.time_rate_diagnostics[i] = glua_tbl_iget_integer(L, i+1, 0); + } + + gk_species.num_time_rate_diagnostics = num_time_rate_diagnostics; + } + with_lua_tbl_tbl(L, "bcs") { int num_bcs = glua_objlen(L); for (int i = 0; i < num_bcs; i++) { @@ -2642,6 +2666,7 @@ gkyl_gyrokinetic_lw_openlibs(lua_State *L) gkyl_register_gyrokinetic_position_map_types(L); gkyl_register_gyrokinetic_collisionless_types(L); gkyl_register_gyrokinetic_field_types(L); + gkyl_register_gyrokinetic_time_rate_diagnostic_types(L); gkyl_register_gyrokinetic_radiation_types(L); gkyl_register_gyrokinetic_radiation_Te_types(L); gkyl_register_gyrokinetic_reaction_types(L); diff --git a/gyrokinetic/apps/gyrokinetic_multib.c b/gyrokinetic/apps/gyrokinetic_multib.c index de0352dee2..365ad66bed 100644 --- a/gyrokinetic/apps/gyrokinetic_multib.c +++ b/gyrokinetic/apps/gyrokinetic_multib.c @@ -175,7 +175,10 @@ singleb_app_new_solver(const struct gkyl_gyrokinetic_multib *mbinp, int bid, for (int n=0; nintegrated_diag_moments[n]; } - species_inp.time_rate_diagnostics = sp->time_rate_diagnostics; + species_inp.num_time_rate_diagnostics = sp->num_time_rate_diagnostics; + for (int n=0; ntime_rate_diagnostics[n]; + } species_inp.boundary_flux_diagnostics = sp->boundary_flux_diagnostics; // Choose proper block-specific species input. diff --git a/gyrokinetic/apps/gyrokinetic_multib_update_ssp_rk3.c b/gyrokinetic/apps/gyrokinetic_multib_update_ssp_rk3.c index 8f079ac882..60aa8e9399 100644 --- a/gyrokinetic/apps/gyrokinetic_multib_update_ssp_rk3.c +++ b/gyrokinetic/apps/gyrokinetic_multib_update_ssp_rk3.c @@ -138,6 +138,7 @@ gyrokinetic_multib_update_ssp_rk3(struct gkyl_gyrokinetic_multib_app* app, doubl // Compute moment of f_old to later compute moment of df/dt. // Do it before the fields are updated, but after dt is calculated. gk_species_calc_int_mom_dt(sbapp, gks, dt, gks->fdot_mom_old); + gk_species_calc_fdot_begin_step(sbapp, gks, dt); } // Compute field energy divided by dt for energy balance diagnostics. @@ -355,6 +356,7 @@ gyrokinetic_multib_update_ssp_rk3(struct gkyl_gyrokinetic_multib_app* app, doubl // Compute moment of f_new to compute moment of df/dt. // Need to do it after the fields are updated. gk_species_calc_int_mom_dt(sbapp, gks, dt, gks->fdot_mom_new); + gk_species_calc_fdot_complete_step(sbapp, gks, dt); } // Compute field energy divided by dt for energy balance diagnostics. @@ -373,4 +375,3 @@ gyrokinetic_multib_update_ssp_rk3(struct gkyl_gyrokinetic_multib_app* app, doubl return st; } - diff --git a/gyrokinetic/apps/gyrokinetic_update_ssp_rk3.c b/gyrokinetic/apps/gyrokinetic_update_ssp_rk3.c index 18a1a7e7a2..e84cbad624 100644 --- a/gyrokinetic/apps/gyrokinetic_update_ssp_rk3.c +++ b/gyrokinetic/apps/gyrokinetic_update_ssp_rk3.c @@ -109,6 +109,7 @@ gyrokinetic_update_ssp_rk3(gkyl_gyrokinetic_app* app, double dt0) // Compute moment of f_old to later compute moment of df/dt. // Do it before the fields are updated, but after dt is calculated. gk_species_calc_int_mom_dt(app, gks, dt, gks->fdot_mom_old); + gk_species_calc_fdot_begin_step(app, gks, dt); } // Compute field energy divided by dt for energy balance diagnostics. @@ -287,6 +288,7 @@ gyrokinetic_update_ssp_rk3(gkyl_gyrokinetic_app* app, double dt0) // Compute moment of f_new to compute moment of df/dt. // Need to do it after the fields are updated. gk_species_calc_int_mom_dt(app, gks, dt, gks->fdot_mom_new); + gk_species_calc_fdot_complete_step(app, gks, dt); } // Scale species according to some criteria. diff --git a/gyrokinetic/creg/rt_gk_asdex_solonly_3x2v_p1.c b/gyrokinetic/creg/rt_gk_asdex_solonly_3x2v_p1.c index 4d14f4918c..7f47d74dce 100644 --- a/gyrokinetic/creg/rt_gk_asdex_solonly_3x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_asdex_solonly_3x2v_p1.c @@ -489,7 +489,11 @@ main(int argc, char **argv) .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, }; // Ion core source: @@ -620,7 +624,11 @@ main(int argc, char **argv) .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, }; // field diff --git a/gyrokinetic/creg/rt_gk_cbc_2x2v_p1.c b/gyrokinetic/creg/rt_gk_cbc_2x2v_p1.c index b49d37b153..74f53f9c76 100644 --- a/gyrokinetic/creg/rt_gk_cbc_2x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_cbc_2x2v_p1.c @@ -764,7 +764,11 @@ main(int argc, char **argv) .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, }; // ions @@ -829,7 +833,11 @@ main(int argc, char **argv) .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, }; // field diff --git a/gyrokinetic/creg/rt_gk_cbc_3x2v_p1.c b/gyrokinetic/creg/rt_gk_cbc_3x2v_p1.c index 2dcfcb6175..fe666cb741 100644 --- a/gyrokinetic/creg/rt_gk_cbc_3x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_cbc_3x2v_p1.c @@ -792,7 +792,11 @@ main(int argc, char **argv) .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, }; // ions @@ -851,7 +855,11 @@ main(int argc, char **argv) .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, }; // field diff --git a/gyrokinetic/creg/rt_gk_d3d_iwl_2x2v_p1.c b/gyrokinetic/creg/rt_gk_d3d_iwl_2x2v_p1.c index 22edb63f91..2037b2f5ac 100644 --- a/gyrokinetic/creg/rt_gk_d3d_iwl_2x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_d3d_iwl_2x2v_p1.c @@ -640,7 +640,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_BIMAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, @@ -730,7 +734,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_BIMAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_d3d_iwl_3x2v_p1.c b/gyrokinetic/creg/rt_gk_d3d_iwl_3x2v_p1.c index 5ac1995229..a66f00434a 100644 --- a/gyrokinetic/creg/rt_gk_d3d_iwl_3x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_d3d_iwl_3x2v_p1.c @@ -722,7 +722,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_MAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, @@ -810,7 +814,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_MAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_helical_zpar_3x2v_p1.c b/gyrokinetic/creg/rt_gk_helical_zpar_3x2v_p1.c index a26038cf30..a8cea39562 100644 --- a/gyrokinetic/creg/rt_gk_helical_zpar_3x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_helical_zpar_3x2v_p1.c @@ -481,7 +481,11 @@ main(int argc, char **argv) .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, }; // ions @@ -557,7 +561,11 @@ main(int argc, char **argv) .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, }; // field diff --git a/gyrokinetic/creg/rt_gk_helical_zvert_3x2v_p1.c b/gyrokinetic/creg/rt_gk_helical_zvert_3x2v_p1.c index b5a364a3f8..611b406dba 100644 --- a/gyrokinetic/creg/rt_gk_helical_zvert_3x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_helical_zvert_3x2v_p1.c @@ -495,7 +495,11 @@ main(int argc, char **argv) .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, }; // ions @@ -571,7 +575,11 @@ main(int argc, char **argv) .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, }; // field diff --git a/gyrokinetic/creg/rt_gk_lapd_cyl_3x2v_p1.c b/gyrokinetic/creg/rt_gk_lapd_cyl_3x2v_p1.c index d98a5c726d..ba9655a0e6 100644 --- a/gyrokinetic/creg/rt_gk_lapd_cyl_3x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_lapd_cyl_3x2v_p1.c @@ -475,7 +475,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_M3PAR, GKYL_F_MOMENT_M3PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, @@ -556,7 +560,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_M3PAR, GKYL_F_MOMENT_M3PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_leaky_bag_1x2v_p1.c b/gyrokinetic/creg/rt_gk_leaky_bag_1x2v_p1.c index 212e976d67..17188d47c9 100644 --- a/gyrokinetic/creg/rt_gk_leaky_bag_1x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_leaky_bag_1x2v_p1.c @@ -310,7 +310,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_MAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_leaky_bag_1x2v_p1_nux_center.c b/gyrokinetic/creg/rt_gk_leaky_bag_1x2v_p1_nux_center.c index 99a5cdde73..00d2ec7549 100644 --- a/gyrokinetic/creg/rt_gk_leaky_bag_1x2v_p1_nux_center.c +++ b/gyrokinetic/creg/rt_gk_leaky_bag_1x2v_p1_nux_center.c @@ -334,7 +334,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_MAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_leaky_bag_1x2v_p1_nux_edge.c b/gyrokinetic/creg/rt_gk_leaky_bag_1x2v_p1_nux_edge.c index ad140fafd0..70b83050e8 100644 --- a/gyrokinetic/creg/rt_gk_leaky_bag_1x2v_p1_nux_edge.c +++ b/gyrokinetic/creg/rt_gk_leaky_bag_1x2v_p1_nux_edge.c @@ -319,7 +319,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_MAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_leaky_bag_1x2v_p1_ux_stretch.c b/gyrokinetic/creg/rt_gk_leaky_bag_1x2v_p1_ux_stretch.c index 4c0061cf68..e4e31bb0f8 100644 --- a/gyrokinetic/creg/rt_gk_leaky_bag_1x2v_p1_ux_stretch.c +++ b/gyrokinetic/creg/rt_gk_leaky_bag_1x2v_p1_ux_stretch.c @@ -310,7 +310,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_MAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_leaky_bag_boltz_sheath_1x2v_p1.c b/gyrokinetic/creg/rt_gk_leaky_bag_boltz_sheath_1x2v_p1.c index 0bfca2a161..2f33e0b5e4 100644 --- a/gyrokinetic/creg/rt_gk_leaky_bag_boltz_sheath_1x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_leaky_bag_boltz_sheath_1x2v_p1.c @@ -243,7 +243,12 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_MAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 3, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_leaky_bag_boltz_sheath_1x2v_p1_nux_center.c b/gyrokinetic/creg/rt_gk_leaky_bag_boltz_sheath_1x2v_p1_nux_center.c index ef8af3c419..5d27a50cb2 100644 --- a/gyrokinetic/creg/rt_gk_leaky_bag_boltz_sheath_1x2v_p1_nux_center.c +++ b/gyrokinetic/creg/rt_gk_leaky_bag_boltz_sheath_1x2v_p1_nux_center.c @@ -251,7 +251,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_MAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_leaky_bag_boltz_sheath_1x2v_p1_nux_edge.c b/gyrokinetic/creg/rt_gk_leaky_bag_boltz_sheath_1x2v_p1_nux_edge.c index eb9e61a90f..72d44819f7 100644 --- a/gyrokinetic/creg/rt_gk_leaky_bag_boltz_sheath_1x2v_p1_nux_edge.c +++ b/gyrokinetic/creg/rt_gk_leaky_bag_boltz_sheath_1x2v_p1_nux_edge.c @@ -250,7 +250,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_MAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_leaky_bag_boltz_sheath_1x2v_p1_ux_stretch.c b/gyrokinetic/creg/rt_gk_leaky_bag_boltz_sheath_1x2v_p1_ux_stretch.c index c2e80b3f94..af331a0751 100644 --- a/gyrokinetic/creg/rt_gk_leaky_bag_boltz_sheath_1x2v_p1_ux_stretch.c +++ b/gyrokinetic/creg/rt_gk_leaky_bag_boltz_sheath_1x2v_p1_ux_stretch.c @@ -242,7 +242,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_MAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_multib_asdex_2x2v_p1.c b/gyrokinetic/creg/rt_gk_multib_asdex_2x2v_p1.c index 7531365a78..62839fd330 100644 --- a/gyrokinetic/creg/rt_gk_multib_asdex_2x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_multib_asdex_2x2v_p1.c @@ -1213,7 +1213,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_BIMAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, @@ -1345,7 +1349,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_BIMAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, diff --git a/gyrokinetic/creg/rt_gk_multib_asdex_solonly_3x2v_p1.c b/gyrokinetic/creg/rt_gk_multib_asdex_solonly_3x2v_p1.c index 01a4d91a74..9b495f313c 100644 --- a/gyrokinetic/creg/rt_gk_multib_asdex_solonly_3x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_multib_asdex_solonly_3x2v_p1.c @@ -718,7 +718,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_BIMAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, @@ -891,7 +895,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_BIMAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, diff --git a/gyrokinetic/creg/rt_gk_multib_nstx_solonly_3x2v_p1.c b/gyrokinetic/creg/rt_gk_multib_nstx_solonly_3x2v_p1.c index c7c81739ce..3bcc60278e 100644 --- a/gyrokinetic/creg/rt_gk_multib_nstx_solonly_3x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_multib_nstx_solonly_3x2v_p1.c @@ -708,7 +708,11 @@ main(int argc, char **argv) }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, .diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, @@ -825,7 +829,11 @@ main(int argc, char **argv) }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, .diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, diff --git a/gyrokinetic/creg/rt_gk_multib_sheath_1x2v_p1.c b/gyrokinetic/creg/rt_gk_multib_sheath_1x2v_p1.c index ecdac989a9..ea66096bf9 100644 --- a/gyrokinetic/creg/rt_gk_multib_sheath_1x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_multib_sheath_1x2v_p1.c @@ -635,7 +635,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_M3PAR, GKYL_F_MOMENT_M3PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, @@ -725,7 +729,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_M3PAR, GKYL_F_MOMENT_M3PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, diff --git a/gyrokinetic/creg/rt_gk_multib_slab_2x2v_p1.c b/gyrokinetic/creg/rt_gk_multib_slab_2x2v_p1.c index e0b2852f62..06f614197a 100644 --- a/gyrokinetic/creg/rt_gk_multib_slab_2x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_multib_slab_2x2v_p1.c @@ -512,7 +512,11 @@ main(int argc, char **argv) }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, @@ -616,7 +620,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_M3PAR, GKYL_F_MOMENT_M3PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, diff --git a/gyrokinetic/creg/rt_gk_multib_step_2x2v_p1.c b/gyrokinetic/creg/rt_gk_multib_step_2x2v_p1.c index 023538ba37..a5db5c242d 100644 --- a/gyrokinetic/creg/rt_gk_multib_step_2x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_multib_step_2x2v_p1.c @@ -1241,7 +1241,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_M3PAR, GKYL_F_MOMENT_M3PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, .diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, @@ -1585,7 +1589,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_M3PAR, GKYL_F_MOMENT_M3PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, .diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, diff --git a/gyrokinetic/creg/rt_gk_multib_step_eirene_2x2v_p1.c b/gyrokinetic/creg/rt_gk_multib_step_eirene_2x2v_p1.c index cd1bb0ba29..f570f932b6 100644 --- a/gyrokinetic/creg/rt_gk_multib_step_eirene_2x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_multib_step_eirene_2x2v_p1.c @@ -1136,7 +1136,11 @@ main(int argc, char **argv) }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, .diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, @@ -1401,7 +1405,11 @@ main(int argc, char **argv) }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, .diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, diff --git a/gyrokinetic/creg/rt_gk_multib_step_nonuniform_2x2v_p1.c b/gyrokinetic/creg/rt_gk_multib_step_nonuniform_2x2v_p1.c index abe718cfdc..5634dd9f24 100644 --- a/gyrokinetic/creg/rt_gk_multib_step_nonuniform_2x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_multib_step_nonuniform_2x2v_p1.c @@ -1360,7 +1360,11 @@ main(int argc, char **argv) }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, .diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, @@ -1732,7 +1736,11 @@ main(int argc, char **argv) }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, .diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, diff --git a/gyrokinetic/creg/rt_gk_multib_step_sol_2x2v_p1.c b/gyrokinetic/creg/rt_gk_multib_step_sol_2x2v_p1.c index 2e2461ab0e..b7cd0fe275 100644 --- a/gyrokinetic/creg/rt_gk_multib_step_sol_2x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_multib_step_sol_2x2v_p1.c @@ -511,7 +511,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_M3PAR, GKYL_F_MOMENT_M3PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, @@ -612,7 +616,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_M3PAR, GKYL_F_MOMENT_M3PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, diff --git a/gyrokinetic/creg/rt_gk_multib_tcv_x21_3x2v_p1.c b/gyrokinetic/creg/rt_gk_multib_tcv_x21_3x2v_p1.c index fd79d97eaf..b54fc5edda 100644 --- a/gyrokinetic/creg/rt_gk_multib_tcv_x21_3x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_multib_tcv_x21_3x2v_p1.c @@ -1005,7 +1005,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_BIMAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, @@ -1130,7 +1134,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_BIMAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, diff --git a/gyrokinetic/creg/rt_gk_neut_sheath_3x2v_p1.c b/gyrokinetic/creg/rt_gk_neut_sheath_3x2v_p1.c index 0f1fe14b3c..7a073a5f04 100644 --- a/gyrokinetic/creg/rt_gk_neut_sheath_3x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_neut_sheath_3x2v_p1.c @@ -733,7 +733,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0M1M2PARM2PERP, GKYL_F_MOMENT_BIMAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, @@ -844,7 +848,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0M1M2PARM2PERP, GKYL_F_MOMENT_BIMAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_sheath_1x2v_p1.c b/gyrokinetic/creg/rt_gk_sheath_1x2v_p1.c index d0f204ba73..993e3577b2 100644 --- a/gyrokinetic/creg/rt_gk_sheath_1x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_sheath_1x2v_p1.c @@ -521,7 +521,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_MAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, @@ -597,7 +601,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_MAXWELLIAN, GKYL_F_MOMENT_BIMAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_sheath_1x2v_p1_ic_import.c b/gyrokinetic/creg/rt_gk_sheath_1x2v_p1_ic_import.c index f7228953db..6997f92a3b 100644 --- a/gyrokinetic/creg/rt_gk_sheath_1x2v_p1_ic_import.c +++ b/gyrokinetic/creg/rt_gk_sheath_1x2v_p1_ic_import.c @@ -522,7 +522,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_MAXWELLIAN, GKYL_F_MOMENT_BIMAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, @@ -599,7 +603,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_MAXWELLIAN, GKYL_F_MOMENT_BIMAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_sheath_2x2v_p1.c b/gyrokinetic/creg/rt_gk_sheath_2x2v_p1.c index 066300afda..f0335d3799 100644 --- a/gyrokinetic/creg/rt_gk_sheath_2x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_sheath_2x2v_p1.c @@ -663,7 +663,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, @@ -750,7 +754,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_sheath_3x2v_p1.c b/gyrokinetic/creg/rt_gk_sheath_3x2v_p1.c index be30c42288..75e5a6af59 100644 --- a/gyrokinetic/creg/rt_gk_sheath_3x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_sheath_3x2v_p1.c @@ -587,7 +587,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, @@ -663,7 +667,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0M1M2PARM2PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_sheath_fluid_neut_1x2v_p1.c b/gyrokinetic/creg/rt_gk_sheath_fluid_neut_1x2v_p1.c index 6ac397d414..10d753fc97 100644 --- a/gyrokinetic/creg/rt_gk_sheath_fluid_neut_1x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_sheath_fluid_neut_1x2v_p1.c @@ -659,7 +659,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_MAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, @@ -769,7 +773,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_MAXWELLIAN }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_sheath_heat_source_2x2v_p1.c b/gyrokinetic/creg/rt_gk_sheath_heat_source_2x2v_p1.c index 024e344de3..8474078ae4 100644 --- a/gyrokinetic/creg/rt_gk_sheath_heat_source_2x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_sheath_heat_source_2x2v_p1.c @@ -770,7 +770,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_M2 }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, @@ -860,7 +864,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_solovev_out_3x2v_p1.c b/gyrokinetic/creg/rt_gk_solovev_out_3x2v_p1.c index 72dc2076c9..abad5d7b0f 100644 --- a/gyrokinetic/creg/rt_gk_solovev_out_3x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_solovev_out_3x2v_p1.c @@ -401,7 +401,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_M3PAR, GKYL_F_MOMENT_M3PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, @@ -480,7 +484,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_M3PAR, GKYL_F_MOMENT_M3PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_step_2x2v_p1_cons.c b/gyrokinetic/creg/rt_gk_step_2x2v_p1_cons.c index 82e3343e00..dd20d03922 100644 --- a/gyrokinetic/creg/rt_gk_step_2x2v_p1_cons.c +++ b/gyrokinetic/creg/rt_gk_step_2x2v_p1_cons.c @@ -426,7 +426,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_M3PAR, GKYL_F_MOMENT_M3PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_M0M1M2PARM2PERP }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, @@ -542,7 +546,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_M3PAR, GKYL_F_MOMENT_M3PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_M0M1M2PARM2PERP }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 0, diff --git a/gyrokinetic/creg/rt_gk_step_out_2x2v_p1.c b/gyrokinetic/creg/rt_gk_step_out_2x2v_p1.c index b9f76cf0c1..1e222adb9c 100644 --- a/gyrokinetic/creg/rt_gk_step_out_2x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_step_out_2x2v_p1.c @@ -398,7 +398,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_M3PAR, GKYL_F_MOMENT_M3PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, @@ -479,7 +483,11 @@ main(int argc, char **argv) .diag_moments = { GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_M3PAR, GKYL_F_MOMENT_M3PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_diag_moments = 0, diff --git a/gyrokinetic/creg/rt_gk_tcv_core_3x2v_p1.c b/gyrokinetic/creg/rt_gk_tcv_core_3x2v_p1.c index 52170f5891..ed35ffc90f 100644 --- a/gyrokinetic/creg/rt_gk_tcv_core_3x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_tcv_core_3x2v_p1.c @@ -636,7 +636,11 @@ main(int argc, char **argv) .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, }; // ions @@ -726,7 +730,11 @@ main(int argc, char **argv) .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, }; // Field. diff --git a/gyrokinetic/creg/rt_gk_tcv_iwl_adapt_source_2x2v_p1.c b/gyrokinetic/creg/rt_gk_tcv_iwl_adapt_source_2x2v_p1.c index 2e0ee0094a..f056066594 100644 --- a/gyrokinetic/creg/rt_gk_tcv_iwl_adapt_source_2x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_tcv_iwl_adapt_source_2x2v_p1.c @@ -665,7 +665,11 @@ main(int argc, char **argv) .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, }; // ions @@ -746,7 +750,11 @@ main(int argc, char **argv) .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, }; struct gkyl_poisson_bias_line target_corner_bcs[] = { diff --git a/gyrokinetic/creg/rt_gk_tcv_iwl_adapt_source_3x2v_p1.c b/gyrokinetic/creg/rt_gk_tcv_iwl_adapt_source_3x2v_p1.c index 1520d5cdde..ae7e88226e 100644 --- a/gyrokinetic/creg/rt_gk_tcv_iwl_adapt_source_3x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_tcv_iwl_adapt_source_3x2v_p1.c @@ -700,7 +700,11 @@ main(int argc, char **argv) .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, }; // ions @@ -785,7 +789,11 @@ main(int argc, char **argv) .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, }; struct gkyl_poisson_bias_line target_corner_bcs[] = { diff --git a/gyrokinetic/creg/rt_gk_wham_1x2v_p1.c b/gyrokinetic/creg/rt_gk_wham_1x2v_p1.c index ac4da07310..71d7a47ca8 100644 --- a/gyrokinetic/creg/rt_gk_wham_1x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_wham_1x2v_p1.c @@ -648,7 +648,11 @@ int main(int argc, char **argv) .diag_moments = {GKYL_F_MOMENT_BIMAXWELLIAN, GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_M3PAR, GKYL_F_MOMENT_M3PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, @@ -737,7 +741,11 @@ int main(int argc, char **argv) .diag_moments = {GKYL_F_MOMENT_BIMAXWELLIAN, GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_M3PAR, GKYL_F_MOMENT_M3PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_HAMILTONIAN }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, diff --git a/gyrokinetic/creg/rt_gk_wham_boltz_elc_poa_1x2v_p1.c b/gyrokinetic/creg/rt_gk_wham_boltz_elc_poa_1x2v_p1.c index cdb60a887f..3b8138195e 100644 --- a/gyrokinetic/creg/rt_gk_wham_boltz_elc_poa_1x2v_p1.c +++ b/gyrokinetic/creg/rt_gk_wham_boltz_elc_poa_1x2v_p1.c @@ -635,7 +635,11 @@ int main(int argc, char **argv) .diag_moments = {GKYL_F_MOMENT_BIMAXWELLIAN, GKYL_F_MOMENT_M0, GKYL_F_MOMENT_M1, GKYL_F_MOMENT_M2, GKYL_F_MOMENT_M2PAR, GKYL_F_MOMENT_M2PERP, GKYL_F_MOMENT_M3PAR, GKYL_F_MOMENT_M3PERP }, .num_integrated_diag_moments = 1, .integrated_diag_moments = { GKYL_F_MOMENT_M0M1M2PARM2PERP }, - .time_rate_diagnostics = true, + .num_time_rate_diagnostics = 2, + .time_rate_diagnostics = { + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, + }, .boundary_flux_diagnostics = { .num_integrated_diag_moments = 1, From 3e97b1f46db9e80933e14613acba1291b1b1e285 Mon Sep 17 00:00:00 2001 From: Maxwell-Rosen Date: Mon, 31 Aug 2026 15:32:40 -0400 Subject: [PATCH 4/7] Refactor Fdot diagnostics handling and streamline write operations in gyrokinetic species --- gyrokinetic/apps/gk_species.c | 106 ++++++++---------- gyrokinetic/apps/gkyl_gyrokinetic_priv.h | 26 ++--- .../apps/gyrokinetic_multib_update_ssp_rk3.c | 35 ++++-- gyrokinetic/apps/gyrokinetic_update_ssp_rk3.c | 28 ++++- 4 files changed, 110 insertions(+), 85 deletions(-) diff --git a/gyrokinetic/apps/gk_species.c b/gyrokinetic/apps/gk_species.c index a0d7490376..1204f534bc 100644 --- a/gyrokinetic/apps/gk_species.c +++ b/gyrokinetic/apps/gk_species.c @@ -349,35 +349,57 @@ gk_species_write_dynamic(gkyl_gyrokinetic_app* app, struct gk_species *gks, doub app->stat.species_io_tm += gkyl_time_diff_now_sec(wst); app->stat.n_io += 1; - if (gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT]) { - struct timespec wdt = gkyl_wall_clock(); + if (!gks->info.is_static && + gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT]) { + // The matching RHS is produced by stage 1 of the next RK step. Queue its + // metadata so it can be written before the forward-Euler update overwrites it. + gks->fdot_io_pending = true; + gks->fdot_io_tm = tm; + gks->fdot_io_frame = frame; + } +} - struct gkyl_msgpack_map_elem io_meta_fdot[] = { - { .key = "Description", .elem_type = GKYL_MP_STRING, - .cval = "Time rate of change of the distribution function." } - }; - int io_meta_fdot_len[] = {gks->io_meta_phase_len, app->gk_geom->io_meta_basic_len, 1}; - const struct gkyl_msgpack_map_elem* io_meta_fdot_union[] = { - gks->io_meta_phase, app->gk_geom->io_meta_basic, io_meta_fdot - }; - struct gkyl_msgpack_data *fdot_mt = gkyl_msgpack_create_union( - sizeof(io_meta_fdot_len)/sizeof(int), io_meta_fdot_len, io_meta_fdot_union); +void +gk_species_write_fdot(gkyl_gyrokinetic_app *app, struct gk_species *gks, + const struct gkyl_array *fdot) +{ + if (!gks->fdot_io_pending) + return; - const char *fdot_fmt = "%s-%s_fdot_%d.gkyl"; - int fdot_sz = gkyl_calc_strlen(fdot_fmt, app->name, gks->info.name, frame); - char fdot_fileNm[fdot_sz+1]; // Ensures no buffer overflow. - snprintf(fdot_fileNm, sizeof fdot_fileNm, fdot_fmt, app->name, gks->info.name, frame); + struct timespec wst = gkyl_wall_clock(); - if (app->use_gpu) { - gkyl_array_copy(gks->fdot_host, gks->fdot); - } - gkyl_comm_array_write(gks->comm, &gks->grid, &gks->local, fdot_mt, - gks->fdot_host, fdot_fileNm); + gkyl_msgpack_map_elem_set_double(gks->io_meta_phase_len, gks->io_meta_phase, + "time", gks->fdot_io_tm); + gkyl_msgpack_map_elem_set_uint(gks->io_meta_phase_len, gks->io_meta_phase, + "frame", gks->fdot_io_frame); - gkyl_msgpack_data_release(fdot_mt); - app->stat.species_diag_io_tm += gkyl_time_diff_now_sec(wdt); - app->stat.n_diag_io += 1; + struct gkyl_msgpack_map_elem io_meta_fdot[] = { + { .key = "Description", .elem_type = GKYL_MP_STRING, + .cval = "Explicit SSP-RK3 stage-1 time rate of change of the distribution function." } + }; + int io_meta_len[] = {gks->io_meta_phase_len, app->gk_geom->io_meta_basic_len, 1}; + const struct gkyl_msgpack_map_elem* io_meta[] = { + gks->io_meta_phase, app->gk_geom->io_meta_basic, io_meta_fdot + }; + struct gkyl_msgpack_data *mt = gkyl_msgpack_create_union( + sizeof(io_meta_len)/sizeof(int), io_meta_len, io_meta); + + const char *fmt = "%s-%s_fdot_%d.gkyl"; + int sz = gkyl_calc_strlen(fmt, app->name, gks->info.name, gks->fdot_io_frame); + char fileNm[sz+1]; // Ensures no buffer overflow. + snprintf(fileNm, sizeof fileNm, fmt, app->name, gks->info.name, gks->fdot_io_frame); + + const struct gkyl_array *fdot_host = fdot; + if (app->use_gpu) { + gkyl_array_copy(gks->f_host, fdot); + fdot_host = gks->f_host; } + gkyl_comm_array_write(gks->comm, &gks->grid, &gks->local, mt, fdot_host, fileNm); + + gks->fdot_io_pending = false; + gkyl_msgpack_data_release(mt); + app->stat.species_diag_io_tm += gkyl_time_diff_now_sec(wst); + app->stat.n_diag_io += 1; } static void @@ -501,26 +523,6 @@ gk_species_calc_int_mom_dt(gkyl_gyrokinetic_app* app, struct gk_species *gks, do gks->calc_int_mom_dt_func(app, gks, dt, fdot_int_mom); } -void -gk_species_calc_fdot_begin_step(gkyl_gyrokinetic_app* app, struct gk_species *gks, double dt) -{ - if (gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT]) { - struct timespec wst = gkyl_wall_clock(); - gkyl_array_set(gks->fdot, -1.0/dt, gks->f); - app->stat.fdot_tm += gkyl_time_diff_now_sec(wst); - } -} - -void -gk_species_calc_fdot_complete_step(gkyl_gyrokinetic_app* app, struct gk_species *gks, double dt) -{ - if (gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT]) { - struct timespec wst = gkyl_wall_clock(); - gkyl_array_accumulate(gks->fdot, 1.0/dt, gks->f); - app->stat.fdot_tm += gkyl_time_diff_now_sec(wst); - } -} - static void gk_species_calc_integrated_mom_dynamic(gkyl_gyrokinetic_app* app, struct gk_species *gks, double tm) { @@ -1505,6 +1507,7 @@ gk_species_init(struct gkyl_gk *gk_app_inp, struct gkyl_gyrokinetic_app *app, st for (int d=0; dtime_rate_diagnostics[d] = false; } + gks->fdot_io_pending = false; for (int d=0; dinfo.num_time_rate_diagnostics; ++d) { enum gkyl_gyrokinetic_time_rate_diagnostic diag = gks->info.time_rate_diagnostics[d]; assert(diag >= 0 && diag < GKYL_GK_TIME_RATE_DIAGNOSTIC_NUM); @@ -1673,14 +1676,6 @@ gk_species_init(struct gkyl_gk *gk_app_inp, struct gkyl_gyrokinetic_app *app, st if (app->use_gpu) gks->f_host = mkarr(false, gks->basis.num_basis, gks->local_ext.volume); - if (gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT]) { - gks->fdot = mkarr(app->use_gpu, gks->basis.num_basis, gks->local_ext.volume); - gks->fdot_host = gks->fdot; - if (app->use_gpu) { - gks->fdot_host = mkarr(false, gks->basis.num_basis, gks->local_ext.volume); - } - } - // Allocate cflrate (scalar array). gks->cflrate = mkarr(app->use_gpu, 1, gks->local_ext.volume); @@ -2152,13 +2147,6 @@ gk_species_release(const gkyl_gyrokinetic_app* app, const struct gk_species *gks gkyl_cu_free(gks->basis_on_dev); } - if (gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT]) { - gkyl_array_release(gks->fdot); - if (app->use_gpu) { - gkyl_array_release(gks->fdot_host); - } - } - gkyl_velocity_map_release(gks->vel_map); gk_species_collisionless_release(app, &gks->collisionless); diff --git a/gyrokinetic/apps/gkyl_gyrokinetic_priv.h b/gyrokinetic/apps/gkyl_gyrokinetic_priv.h index 076736cb43..f9faa60161 100644 --- a/gyrokinetic/apps/gkyl_gyrokinetic_priv.h +++ b/gyrokinetic/apps/gkyl_gyrokinetic_priv.h @@ -1030,7 +1030,6 @@ struct gk_species { int io_meta_conf_len; // Number of elements in io_meta_conf. struct gkyl_array *f, *f1, *fnew; // Arrays for updates. - struct gkyl_array *fdot, *fdot_host; // Phase-space (f_new-f_old)/dt and its host copy. struct gkyl_array *cflrate; // CFL rate in each cell. struct gkyl_array *cflrate_ho; // CFL rate in each cell on host-side. @@ -1056,6 +1055,9 @@ struct gk_species { // Lookup table populated from info.time_rate_diagnostics. bool time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_NUM]; + bool fdot_io_pending; // Whether the next stage-1 RHS should be written. + double fdot_io_tm; // Time associated with the pending Fdot frame. + int fdot_io_frame; // Frame associated with the pending Fdot write. struct gkyl_array_integrate* integ_wfsq_op; // Operator to integrate w*f^2. double *L2norm_local, *L2norm_global; // L2norm in local MPI process and across the communicator. @@ -3282,6 +3284,16 @@ void gk_species_n_iter_corr(gkyl_gyrokinetic_app *app, const struct gk_species * */ void gk_species_write(gkyl_gyrokinetic_app* app, struct gk_species *gks, double tm, int frame); +/** + * Write a queued Fdot diagnostic from an SSP-RK3 stage-1 RHS. + * + * @param app Gyrokinetic app object. + * @param gks Species object. + * @param fdot Stage-1 RHS to write. + */ +void gk_species_write_fdot(gkyl_gyrokinetic_app *app, struct gk_species *gks, + const struct gkyl_array *fdot); + /** * Species moment write function. * @@ -3338,18 +3350,6 @@ void gk_species_write_L2norm(gkyl_gyrokinetic_app* app, struct gk_species *gks); void gk_species_calc_int_mom_dt(gkyl_gyrokinetic_app* app, struct gk_species *gks, double dt, struct gkyl_array *fdot_int_mom); -/** - * Store the old-distribution contribution to the phase-space fdot diagnostic. - */ -void -gk_species_calc_fdot_begin_step(gkyl_gyrokinetic_app* app, struct gk_species *gks, double dt); - -/** - * Add the new-distribution contribution to the phase-space fdot diagnostic. - */ -void -gk_species_calc_fdot_complete_step(gkyl_gyrokinetic_app* app, struct gk_species *gks, double dt); - /** * Delete resources used in species. * diff --git a/gyrokinetic/apps/gyrokinetic_multib_update_ssp_rk3.c b/gyrokinetic/apps/gyrokinetic_multib_update_ssp_rk3.c index 60aa8e9399..7c14964e6b 100644 --- a/gyrokinetic/apps/gyrokinetic_multib_update_ssp_rk3.c +++ b/gyrokinetic/apps/gyrokinetic_multib_update_ssp_rk3.c @@ -1,12 +1,34 @@ #include +typedef void (*write_fdot_func_t)(struct gkyl_gyrokinetic_multib_app *app, + struct gkyl_array *fout[]); + +static void +write_fdot_enabled(struct gkyl_gyrokinetic_multib_app *app, + struct gkyl_array *fout[]) +{ + for (int b=0; bnum_local_blocks; ++b) { + struct gkyl_gyrokinetic_app *sbapp = app->singleb_apps[b]; + int li_charged = b * app->num_species; + for (int i=0; inum_species; ++i) + gk_species_write_fdot(sbapp, &sbapp->species[i], fout[li_charged+i]); + } +} + +static void +write_fdot_disabled(struct gkyl_gyrokinetic_multib_app *app, + struct gkyl_array *fout[]) +{ + // Do nothing. +} + static void gyrokinetic_multib_forward_euler(struct gkyl_gyrokinetic_multib_app* app, double tcurr, double dt, const struct gkyl_array *fin[], struct gkyl_array *fout[], struct gkyl_array **bflux_in[], struct gkyl_array **bflux_out[], const struct gkyl_array *fin_neut[], struct gkyl_array *fout_neut[], struct gkyl_array **bflux_in_neut[], struct gkyl_array **bflux_out_neut[], - struct gkyl_update_status *st) + write_fdot_func_t write_fdot_func, struct gkyl_update_status *st) { struct timespec wst_fe = gkyl_wall_clock(); // Take a forward Euler step with the suggested time-step dt. This may @@ -34,6 +56,8 @@ gyrokinetic_multib_forward_euler(struct gkyl_gyrokinetic_multib_app* app, double st->dt_actual = dtmin_global; app->stat.dfdt_dt_reduce_tm += gkyl_time_diff_now_sec(wtm); + write_fdot_func(app, fout); + struct timespec wst = gkyl_wall_clock(); // Complete update of distribution functions. double dta = st->dt_actual; @@ -109,7 +133,7 @@ gyrokinetic_multib_update_ssp_rk3(struct gkyl_gyrokinetic_multib_app* app, doubl } gyrokinetic_multib_forward_euler(app, tcurr, dt, fin, fout, bflux_in, bflux_out, - fin_neut, fout_neut, bflux_in_neut, bflux_out_neut, &st); + fin_neut, fout_neut, bflux_in_neut, bflux_out_neut, write_fdot_enabled, &st); dt = st.dt_actual; // Subtract boundary flux f from f1 so that we only step boundary @@ -138,7 +162,6 @@ gyrokinetic_multib_update_ssp_rk3(struct gkyl_gyrokinetic_multib_app* app, doubl // Compute moment of f_old to later compute moment of df/dt. // Do it before the fields are updated, but after dt is calculated. gk_species_calc_int_mom_dt(sbapp, gks, dt, gks->fdot_mom_old); - gk_species_calc_fdot_begin_step(sbapp, gks, dt); } // Compute field energy divided by dt for energy balance diagnostics. @@ -175,7 +198,7 @@ gyrokinetic_multib_update_ssp_rk3(struct gkyl_gyrokinetic_multib_app* app, doubl } gyrokinetic_multib_forward_euler(app, tcurr+dt, dt, fin, fout, bflux_in, bflux_out, - fin_neut, fout_neut, bflux_in_neut, bflux_out_neut, &st); + fin_neut, fout_neut, bflux_in_neut, bflux_out_neut, write_fdot_disabled, &st); if (st.dt_actual < dt) { @@ -264,7 +287,7 @@ gyrokinetic_multib_update_ssp_rk3(struct gkyl_gyrokinetic_multib_app* app, doubl } gyrokinetic_multib_forward_euler(app, tcurr+dt/2, dt, fin, fout, bflux_in, bflux_out, - fin_neut, fout_neut, bflux_in_neut, bflux_out_neut, &st); + fin_neut, fout_neut, bflux_in_neut, bflux_out_neut, write_fdot_disabled, &st); if (st.dt_actual < dt) { // Recalculate the field. @@ -356,7 +379,6 @@ gyrokinetic_multib_update_ssp_rk3(struct gkyl_gyrokinetic_multib_app* app, doubl // Compute moment of f_new to compute moment of df/dt. // Need to do it after the fields are updated. gk_species_calc_int_mom_dt(sbapp, gks, dt, gks->fdot_mom_new); - gk_species_calc_fdot_complete_step(sbapp, gks, dt); } // Compute field energy divided by dt for energy balance diagnostics. @@ -374,4 +396,3 @@ gyrokinetic_multib_update_ssp_rk3(struct gkyl_gyrokinetic_multib_app* app, doubl return st; } - diff --git a/gyrokinetic/apps/gyrokinetic_update_ssp_rk3.c b/gyrokinetic/apps/gyrokinetic_update_ssp_rk3.c index e84cbad624..747d9bd9bc 100644 --- a/gyrokinetic/apps/gyrokinetic_update_ssp_rk3.c +++ b/gyrokinetic/apps/gyrokinetic_update_ssp_rk3.c @@ -1,12 +1,28 @@ #include +typedef void (*write_fdot_func_t)(gkyl_gyrokinetic_app *app, + struct gkyl_array *fout[]); + +static void +write_fdot_enabled(gkyl_gyrokinetic_app *app, struct gkyl_array *fout[]) +{ + for (int i=0; inum_species; ++i) + gk_species_write_fdot(app, &app->species[i], fout[i]); +} + +static void +write_fdot_disabled(gkyl_gyrokinetic_app *app, struct gkyl_array *fout[]) +{ + // Do nothing. +} + static void gyrokinetic_forward_euler(gkyl_gyrokinetic_app* app, double tcurr, double dt, const struct gkyl_array *fin[], struct gkyl_array *fout[], struct gkyl_array **bflux_in[], struct gkyl_array **bflux_out[], const struct gkyl_array *fin_neut[], struct gkyl_array *fout_neut[], struct gkyl_array **bflux_in_neut[], struct gkyl_array **bflux_out_neut[], - struct gkyl_update_status *st) + write_fdot_func_t write_fdot_func, struct gkyl_update_status *st) { struct timespec wst_fe = gkyl_wall_clock(); @@ -20,6 +36,8 @@ gyrokinetic_forward_euler(gkyl_gyrokinetic_app* app, double tcurr, double dt, // Compute the time rate of change of the distributions, df/dt. gyrokinetic_rhs(app, tcurr, dt, fin, fout, bflux_out, fin_neut, fout_neut, bflux_out_neut, st); + write_fdot_func(app, fout); + struct timespec wst = gkyl_wall_clock(); // Complete update of distribution functions. double dta = st->dt_actual; @@ -87,7 +105,7 @@ gyrokinetic_update_ssp_rk3(gkyl_gyrokinetic_app* app, double dt0) } gyrokinetic_forward_euler(app, tcurr, dt, fin, fout, bflux_in, bflux_out, - fin_neut, fout_neut, bflux_in_neut, bflux_out_neut, &st); + fin_neut, fout_neut, bflux_in_neut, bflux_out_neut, write_fdot_enabled, &st); dt = st.dt_actual; // Subtract boundary flux f from f1 so that we only step boundary @@ -109,7 +127,6 @@ gyrokinetic_update_ssp_rk3(gkyl_gyrokinetic_app* app, double dt0) // Compute moment of f_old to later compute moment of df/dt. // Do it before the fields are updated, but after dt is calculated. gk_species_calc_int_mom_dt(app, gks, dt, gks->fdot_mom_old); - gk_species_calc_fdot_begin_step(app, gks, dt); } // Compute field energy divided by dt for energy balance diagnostics. @@ -140,7 +157,7 @@ gyrokinetic_update_ssp_rk3(gkyl_gyrokinetic_app* app, double dt0) } gyrokinetic_forward_euler(app, tcurr+dt, dt, fin, fout, bflux_in, bflux_out, - fin_neut, fout_neut, bflux_in_neut, bflux_out_neut, &st); + fin_neut, fout_neut, bflux_in_neut, bflux_out_neut, write_fdot_disabled, &st); if (st.dt_actual < dt) { @@ -212,7 +229,7 @@ gyrokinetic_update_ssp_rk3(gkyl_gyrokinetic_app* app, double dt0) } gyrokinetic_forward_euler(app, tcurr+dt/2, dt, fin, fout, bflux_in, bflux_out, - fin_neut, fout_neut, bflux_in_neut, bflux_out_neut, &st); + fin_neut, fout_neut, bflux_in_neut, bflux_out_neut, write_fdot_disabled, &st); if (st.dt_actual < dt) { // Recalculate the field. @@ -288,7 +305,6 @@ gyrokinetic_update_ssp_rk3(gkyl_gyrokinetic_app* app, double dt0) // Compute moment of f_new to compute moment of df/dt. // Need to do it after the fields are updated. gk_species_calc_int_mom_dt(app, gks, dt, gks->fdot_mom_new); - gk_species_calc_fdot_complete_step(app, gks, dt); } // Scale species according to some criteria. From b6783f801afc9e1d8b2506f1502ba8dcfe1d6fe7 Mon Sep 17 00:00:00 2001 From: Maxwell-Rosen Date: Mon, 31 Aug 2026 15:34:46 -0400 Subject: [PATCH 5/7] Clarify comment for time-rate diagnostic enum in gkyl_gyrokinetic.h --- gyrokinetic/apps/gkyl_gyrokinetic.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/gyrokinetic/apps/gkyl_gyrokinetic.h b/gyrokinetic/apps/gkyl_gyrokinetic.h index cf8ddff4cf..d54c5428ec 100644 --- a/gyrokinetic/apps/gkyl_gyrokinetic.h +++ b/gyrokinetic/apps/gkyl_gyrokinetic.h @@ -100,7 +100,7 @@ enum gkyl_gyrokinetic_time_rate_diagnostic { GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT = 0, // Phase-space (f_new-f_old)/dt. GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, // Volume-integrated moments of fdot. GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, // Volume integral of the absolute fdot moments. - GKYL_GK_TIME_RATE_DIAGNOSTIC_NUM, + GKYL_GK_TIME_RATE_DIAGNOSTIC_NUM, // Sentinal value. Not a diagnostic. Index for total number of time-rate diagnostics. }; // Parameters for collisionless terms. From 51f2502ebaa06c83fb66f2723d33e57bd5b42778 Mon Sep 17 00:00:00 2001 From: Maxwell-Rosen Date: Tue, 1 Sep 2026 11:51:04 -0400 Subject: [PATCH 6/7] I noticed that this PR implemented gkyl_array_reduce_range_sum_abs as a seperate function rather than putting it inside gkyl_array_reduce_range as an option for GKYL_SUM_ABS. The existing implementation is correct, so refactor this to use the gkyl_array_reduce_range method --- core/unit/ctest_array_reduce.c | 4 +-- core/zero/array_dg_reduce.c | 13 ++++++++- core/zero/array_reduce.c | 43 ++++++++++++----------------- core/zero/gkyl_array_reduce.h | 10 ------- core/zero/gkyl_elem_type.h | 2 +- gkeyll/lua/DataStruct/ZeroArray.lua | 4 ++- gyrokinetic/apps/gk_species.c | 2 +- 7 files changed, 37 insertions(+), 41 deletions(-) diff --git a/core/unit/ctest_array_reduce.c b/core/unit/ctest_array_reduce.c index 60bd98f341..954b1f931a 100644 --- a/core/unit/ctest_array_reduce.c +++ b/core/unit/ctest_array_reduce.c @@ -116,7 +116,7 @@ void test_sum_abs_reduce_range() } double asum[2]; - gkyl_array_reduce_range_sum_abs(asum, arr, &range); + gkyl_array_reduce_range(asum, arr, GKYL_SUM_ABS, &range); TEST_CHECK( asum[0] == 0.5*range.volume ); TEST_CHECK( asum[1] == 1.5*range.volume ); @@ -144,7 +144,7 @@ void test_cu_sum_abs_reduce_range() double asum[2]; double *asum_cu = gkyl_cu_malloc(2*sizeof(double)); - gkyl_array_reduce_range_sum_abs(asum_cu, arr_cu, &range); + gkyl_array_reduce_range(asum_cu, arr_cu, GKYL_SUM_ABS, &range); gkyl_cu_memcpy(asum, asum_cu, sizeof(asum), GKYL_CU_MEMCPY_D2H); TEST_CHECK( asum[0] == 0.5*range.volume ); diff --git a/core/zero/array_dg_reduce.c b/core/zero/array_dg_reduce.c index 0414f3d424..7543e449d6 100644 --- a/core/zero/array_dg_reduce.c +++ b/core/zero/array_dg_reduce.c @@ -21,6 +21,9 @@ gkyl_array_dg_reducec(double *out, const struct gkyl_array *arr, int comp, case GKYL_SUM: gkyl_array_dg_reducec_sum_cu(out, arr, comp, basis); break; + default: + assert(false); + break; } return; } @@ -66,6 +69,9 @@ gkyl_array_dg_reducec(double *out, const struct gkyl_array *arr, int comp, } } break; + default: + assert(false); + break; } } @@ -87,6 +93,9 @@ gkyl_array_dg_reducec_range(double *out, const struct gkyl_array *arr, int comp, case GKYL_SUM: gkyl_array_dg_reducec_range_sum_cu(out, arr, comp, basis, range); break; + default: + assert(false); + break; } return; } @@ -138,6 +147,8 @@ gkyl_array_dg_reducec_range(double *out, const struct gkyl_array *arr, int comp, } } break; + default: + assert(false); + break; } } - diff --git a/core/zero/array_reduce.c b/core/zero/array_reduce.c index 5890711f53..c300de1acc 100644 --- a/core/zero/array_reduce.c +++ b/core/zero/array_reduce.c @@ -20,6 +20,9 @@ gkyl_array_reduce(double *out, const struct gkyl_array *arr, enum gkyl_array_op case GKYL_SUM: gkyl_array_reduce_sum_cu(out, arr); break; + default: + assert(false); + break; } return; } @@ -55,6 +58,9 @@ gkyl_array_reduce(double *out, const struct gkyl_array *arr, enum gkyl_array_op out[k] += d[k]; } break; + default: + assert(false); + break; } } @@ -76,6 +82,9 @@ gkyl_array_reduce_range(double *res, case GKYL_SUM: gkyl_array_reduce_range_sum_cu(res, arr, range); break; + case GKYL_SUM_ABS: + gkyl_array_reduce_range_sum_abs_cu(res, arr, range); + break; } return; } @@ -116,31 +125,15 @@ gkyl_array_reduce_range(double *res, res[i] += d[i]; } break; - } -} - -void -gkyl_array_reduce_range_sum_abs(double *res, - const struct gkyl_array *arr, const struct gkyl_range *range) -{ - assert(arr->type == GKYL_DOUBLE); - -#ifdef GKYL_HAVE_CUDA - if (gkyl_array_is_cu_dev(arr)) { - gkyl_array_reduce_range_sum_abs_cu(res, arr, range); - return; - } -#endif - - long n = arr->ncomp; - for (long i=0; ired_integ_diag, gks->fdot_mom_new, &app->local); + gkyl_array_reduce_range(gks->red_integ_diag, gks->fdot_mom_new, GKYL_SUM_ABS, &app->local); gkyl_comm_allreduce(app->comm, GKYL_DOUBLE, GKYL_SUM, num_mom, gks->red_integ_diag, gks->red_integ_diag_global); if (app->use_gpu) { From 6d8b4573d32f1efb6109e706b6b530c32451f514 Mon Sep 17 00:00:00 2001 From: Maxwell-Rosen Date: Thu, 3 Sep 2026 10:44:36 -0400 Subject: [PATCH 7/7] Since the integrated fdot and fdot_abs is computing moments of fdot, add diagnostics that write the moments of fdot and fdot_abs that are computed from Implemented FdotMoments and FdotAbsMoments diagnostics. - Reuses the existing RK-computed moment arrays without recalculation. - Writes P0 configuration-space files named *_fdot__.gkyl and *_fdot_abs__.gkyl. - Added C enums and Lua bindings in gyrokinetic/apps/gkyl_gyrokinetic.h:103 and gyrokinetic/apps/gyrokinetic_lw.c:83. - CPU regression verified absolute values and agreement with both integrated diagnostics. - Library compiled and linked successfully. The final build step only failed when copying shared data outside the writable workspace. Avoid if statements in time loops - Time-loop fdot calculation now uses preselected enabled/no-op function pointers. - Signed, absolute, combined, and disabled writers are selected during initialization. - All new conditionals are limited to initialization/release. - Enabled/disabled regressions and Valgrind passed. - Library compiled and linked; only the sandboxed post-build data copy failed. --- gyrokinetic/apps/gk_species.c | 160 +++++++++++++++++- gyrokinetic/apps/gkyl_gyrokinetic.h | 2 + gyrokinetic/apps/gkyl_gyrokinetic_priv.h | 11 ++ gyrokinetic/apps/gyrokinetic_lw.c | 2 + .../apps/gyrokinetic_multib_update_ssp_rk3.c | 1 + gyrokinetic/apps/gyrokinetic_update_ssp_rk3.c | 1 + 6 files changed, 171 insertions(+), 6 deletions(-) diff --git a/gyrokinetic/apps/gk_species.c b/gyrokinetic/apps/gk_species.c index 3ab08e4381..7fa0873296 100644 --- a/gyrokinetic/apps/gk_species.c +++ b/gyrokinetic/apps/gk_species.c @@ -501,6 +501,113 @@ gk_species_write_mom_static(gkyl_gyrokinetic_app* app, struct gk_species *gks, d // do nothing } +static bool +gk_species_has_fdot_mom_diagnostic(const struct gk_species *gks) +{ + return !gks->info.is_static && ( + gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_MOMENTS] || + gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_MOMENTS] || + gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS] || + gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS]); +} + +static void +gk_species_write_one_fdot_mom(gkyl_gyrokinetic_app* app, struct gk_species *gks, + double tm, int frame, const char *fdot_name, const char *description) +{ + struct timespec wst = gkyl_wall_clock(); + + gkyl_msgpack_map_elem_set_double(gks->io_meta_conf_len, gks->io_meta_conf, "time", tm); + gkyl_msgpack_map_elem_set_uint(gks->io_meta_conf_len, gks->io_meta_conf, "frame", frame); + + enum gkyl_distribution_moments mom_type = gks->info.num_integrated_diag_moments == 0 ? + GKYL_F_MOMENT_M0M1M2PARM2PERP : gks->info.integrated_diag_moments[0]; + const char *mom_name = gkyl_distribution_moments_strs[mom_type]; + + struct gkyl_msgpack_map_elem io_meta_fdot_mom[] = { + { .key = "poly_order", .elem_type = GKYL_MP_UNSIGNED_INT, .uval = 0 }, + { .key = "basis_type", .elem_type = GKYL_MP_STRING, .cval = "serendipity" }, + { .key = "Description", .elem_type = GKYL_MP_STRING, .cval = (char*)description } + }; + int io_meta_len[] = { + gks->io_meta_conf_len, app->gk_geom->io_meta_basic_len, + sizeof(io_meta_fdot_mom)/sizeof(io_meta_fdot_mom[0]) + }; + const struct gkyl_msgpack_map_elem* io_meta[] = { + gks->io_meta_conf, app->gk_geom->io_meta_basic, io_meta_fdot_mom + }; + struct gkyl_msgpack_data *mt = gkyl_msgpack_create_union( + sizeof(io_meta_len)/sizeof(int), io_meta_len, io_meta); + + const char *fmt = "%s-%s_%s_%s_%d.gkyl"; + int sz = gkyl_calc_strlen(fmt, app->name, gks->info.name, fdot_name, mom_name, frame); + char fileNm[sz+1]; // Ensures no buffer overflow. + snprintf(fileNm, sizeof fileNm, fmt, app->name, gks->info.name, fdot_name, mom_name, frame); + + gkyl_comm_array_write(app->comm, &app->grid, &app->local, mt, + gks->integ_moms.marr_host, fileNm); + + gkyl_msgpack_data_release(mt); + app->stat.species_diag_io_tm += gkyl_time_diff_now_sec(wst); + app->stat.n_diag_io += 1; +} + +static void +gk_species_write_fdot_mom_disabled(gkyl_gyrokinetic_app* app, struct gk_species *gks, + double tm, int frame) +{ +} + +static void +gk_species_copy_fdot_mom_to_host(struct gk_species *gks) +{ + gkyl_array_copy(gks->integ_moms.marr_host, gks->fdot_mom_new); +} + +static void +gk_species_abs_fdot_mom_host(gkyl_gyrokinetic_app* app, struct gk_species *gks) +{ + struct gkyl_range_iter iter; + gkyl_range_iter_init(&iter, &app->local); + while (gkyl_range_iter_next(&iter)) { + long loc = gkyl_range_idx(&app->local, iter.idx); + double *fdot_mom = gkyl_array_fetch(gks->integ_moms.marr_host, loc); + for (int m=0; minteg_moms.num_mom; ++m) + fdot_mom[m] = fabs(fdot_mom[m]); + } +} + +static void +gk_species_write_fdot_mom_enabled(gkyl_gyrokinetic_app* app, struct gk_species *gks, + double tm, int frame) +{ + gk_species_copy_fdot_mom_to_host(gks); + gk_species_write_one_fdot_mom(app, gks, tm, frame, "fdot", + "Configuration-space moments of (f_new-f_old)/dt."); +} + +static void +gk_species_write_fdot_abs_mom_enabled(gkyl_gyrokinetic_app* app, struct gk_species *gks, + double tm, int frame) +{ + gk_species_copy_fdot_mom_to_host(gks); + gk_species_abs_fdot_mom_host(app, gks); + gk_species_write_one_fdot_mom(app, gks, tm, frame, "fdot_abs", + "Absolute value of the configuration-space moments of (f_new-f_old)/dt."); +} + +static void +gk_species_write_fdot_and_abs_mom_enabled(gkyl_gyrokinetic_app* app, struct gk_species *gks, + double tm, int frame) +{ + gk_species_copy_fdot_mom_to_host(gks); + gk_species_write_one_fdot_mom(app, gks, tm, frame, "fdot", + "Configuration-space moments of (f_new-f_old)/dt."); + gk_species_abs_fdot_mom_host(app, gks); + gk_species_write_one_fdot_mom(app, gks, tm, frame, "fdot_abs", + "Absolute value of the configuration-space moments of (f_new-f_old)/dt."); +} + static void gk_species_calc_int_mom_dt_enabled(gkyl_gyrokinetic_app* app, struct gk_species *gks, double dt, struct gkyl_array *fdot_int_mom) { @@ -523,6 +630,26 @@ gk_species_calc_int_mom_dt(gkyl_gyrokinetic_app* app, struct gk_species *gks, do gks->calc_int_mom_dt_func(app, gks, dt, fdot_int_mom); } +static void +gk_species_calc_fdot_mom_enabled(gkyl_gyrokinetic_app* app, struct gk_species *gks) +{ + struct timespec wst = gkyl_wall_clock(); + gkyl_array_accumulate(gks->fdot_mom_new, -1.0, gks->fdot_mom_old); + gks->write_fdot_mom_func = gks->write_fdot_mom_ready_func; + app->stat.fdot_tm += gkyl_time_diff_now_sec(wst); +} + +static void +gk_species_calc_fdot_mom_disabled(gkyl_gyrokinetic_app* app, struct gk_species *gks) +{ +} + +void +gk_species_calc_fdot_mom(gkyl_gyrokinetic_app* app, struct gk_species *gks) +{ + gks->calc_fdot_mom_func(app, gks); +} + static void gk_species_calc_integrated_mom_dynamic(gkyl_gyrokinetic_app* app, struct gk_species *gks, double tm) { @@ -551,8 +678,6 @@ gk_species_calc_integrated_mom_dynamic(gkyl_gyrokinetic_app* app, struct gk_spec bool calc_fdot_abs_integ = gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS]; if (calc_fdot_integ || calc_fdot_abs_integ) { - gkyl_array_accumulate(gks->fdot_mom_new, -1.0, gks->fdot_mom_old); - if (calc_fdot_integ) { // Sum the signed time-rate moment over the whole domain. gkyl_array_reduce_range(gks->red_integ_diag, gks->fdot_mom_new, GKYL_SUM, &app->local); @@ -849,7 +974,7 @@ gk_species_release_dynamic(const gkyl_gyrokinetic_app* app, const struct gk_spec s->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS]; bool calc_fdot_abs_integ = s->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS]; - if (calc_fdot_integ || calc_fdot_abs_integ) { + if (gk_species_has_fdot_mom_diagnostic(s)) { gkyl_array_release(s->fdot_mom_old); gkyl_array_release(s->fdot_mom_new); } @@ -920,9 +1045,15 @@ gk_species_init_dynamic(struct gkyl_gk *gk_app_inp, struct gkyl_gyrokinetic_app gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS]; bool calc_fdot_abs_integ = gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS]; - if (calc_fdot_integ || calc_fdot_abs_integ) { + bool write_fdot_mom = + gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_MOMENTS]; + bool write_fdot_abs_mom = + gks->time_rate_diagnostics[GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_MOMENTS]; + if (gk_species_has_fdot_mom_diagnostic(gks)) { gks->fdot_mom_old = mkarr(app->use_gpu, gks->integ_moms.marr->ncomp, gks->integ_moms.marr->size); gks->fdot_mom_new = mkarr(app->use_gpu, gks->integ_moms.marr->ncomp, gks->integ_moms.marr->size); + gkyl_array_clear(gks->fdot_mom_old, 0.0); + gkyl_array_clear(gks->fdot_mom_new, 0.0); } if (calc_fdot_integ) { gks->fdot_integ_diag = gkyl_dynvec_new(GKYL_DOUBLE, gks->integ_moms.num_mom); @@ -1175,10 +1306,23 @@ gk_species_init_dynamic(struct gkyl_gk *gk_app_inp, struct gkyl_gyrokinetic_app gks->write_integrated_mom_func = gk_species_write_integrated_mom_dynamic; gks->calc_L2norm_func = gk_species_calc_L2norm_dynamic; gks->write_L2norm_func = gk_species_write_L2norm_dynamic; - if (calc_fdot_integ || calc_fdot_abs_integ) + if (gk_species_has_fdot_mom_diagnostic(gks)) { gks->calc_int_mom_dt_func = gk_species_calc_int_mom_dt_enabled; - else + gks->calc_fdot_mom_func = gk_species_calc_fdot_mom_enabled; + } + else { gks->calc_int_mom_dt_func = gk_species_calc_int_mom_dt_disabled; + gks->calc_fdot_mom_func = gk_species_calc_fdot_mom_disabled; + } + gks->write_fdot_mom_func = gk_species_write_fdot_mom_disabled; + if (write_fdot_mom && write_fdot_abs_mom) + gks->write_fdot_mom_ready_func = gk_species_write_fdot_and_abs_mom_enabled; + else if (write_fdot_mom) + gks->write_fdot_mom_ready_func = gk_species_write_fdot_mom_enabled; + else if (write_fdot_abs_mom) + gks->write_fdot_mom_ready_func = gk_species_write_fdot_abs_mom_enabled; + else + gks->write_fdot_mom_ready_func = gk_species_write_fdot_mom_disabled; } // Initialize static species object. @@ -1205,6 +1349,9 @@ gk_species_init_static(struct gkyl_gk *gk_app_inp, struct gkyl_gyrokinetic_app * gks->calc_L2norm_func = gk_species_calc_L2norm_static; gks->write_L2norm_func = gk_species_write_L2norm_static; gks->calc_int_mom_dt_func = gk_species_calc_int_mom_dt_disabled; + gks->calc_fdot_mom_func = gk_species_calc_fdot_mom_disabled; + gks->write_fdot_mom_func = gk_species_write_fdot_mom_disabled; + gks->write_fdot_mom_ready_func = gk_species_write_fdot_mom_disabled; } void @@ -2098,6 +2245,7 @@ gk_species_write_mom(gkyl_gyrokinetic_app* app, struct gk_species *gks, double t } else gks->write_mom_func(app, gks, tm, frame); + gks->write_fdot_mom_func(app, gks, tm, frame); } void diff --git a/gyrokinetic/apps/gkyl_gyrokinetic.h b/gyrokinetic/apps/gkyl_gyrokinetic.h index d54c5428ec..4cded34be0 100644 --- a/gyrokinetic/apps/gkyl_gyrokinetic.h +++ b/gyrokinetic/apps/gkyl_gyrokinetic.h @@ -100,6 +100,8 @@ enum gkyl_gyrokinetic_time_rate_diagnostic { GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT = 0, // Phase-space (f_new-f_old)/dt. GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS, // Volume-integrated moments of fdot. GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS, // Volume integral of the absolute fdot moments. + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_MOMENTS, // Configuration-space moments of fdot. + GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_MOMENTS, // Absolute value of the configuration-space fdot moments. GKYL_GK_TIME_RATE_DIAGNOSTIC_NUM, // Sentinal value. Not a diagnostic. Index for total number of time-rate diagnostics. }; diff --git a/gyrokinetic/apps/gkyl_gyrokinetic_priv.h b/gyrokinetic/apps/gkyl_gyrokinetic_priv.h index f9faa60161..b5acd7fba9 100644 --- a/gyrokinetic/apps/gkyl_gyrokinetic_priv.h +++ b/gyrokinetic/apps/gkyl_gyrokinetic_priv.h @@ -1151,6 +1151,9 @@ struct gk_species { void (*calc_L2norm_func)(gkyl_gyrokinetic_app* app, struct gk_species *gks, double tm); void (*write_L2norm_func)(gkyl_gyrokinetic_app* app, struct gk_species *gks); void (*calc_int_mom_dt_func)(gkyl_gyrokinetic_app* app, struct gk_species *gks, double dt, struct gkyl_array *fdot_int_mom); + void (*calc_fdot_mom_func)(gkyl_gyrokinetic_app* app, struct gk_species *gks); + void (*write_fdot_mom_func)(gkyl_gyrokinetic_app* app, struct gk_species *gks, double tm, int frame); + void (*write_fdot_mom_ready_func)(gkyl_gyrokinetic_app* app, struct gk_species *gks, double tm, int frame); // Quantities used for FLR model: struct gkyl_array *m0_gyroavg; // Gyroaveraged particle density. @@ -3350,6 +3353,14 @@ void gk_species_write_L2norm(gkyl_gyrokinetic_app* app, struct gk_species *gks); void gk_species_calc_int_mom_dt(gkyl_gyrokinetic_app* app, struct gk_species *gks, double dt, struct gkyl_array *fdot_int_mom); +/** Finish computing the finite-difference fdot moments after a time step. + * + * @param app Gyrokinetic app object. + * @param gks Gyrokinetic species object. + */ +void +gk_species_calc_fdot_mom(gkyl_gyrokinetic_app* app, struct gk_species *gks); + /** * Delete resources used in species. * diff --git a/gyrokinetic/apps/gyrokinetic_lw.c b/gyrokinetic/apps/gyrokinetic_lw.c index 4b3e2b7a4f..81f778fae5 100644 --- a/gyrokinetic/apps/gyrokinetic_lw.c +++ b/gyrokinetic/apps/gyrokinetic_lw.c @@ -80,6 +80,8 @@ static const struct gkyl_str_int_pair gk_time_rate_diagnostic_type[] = { { "Fdot", GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT }, { "FdotIntegratedMoments", GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_INTEGRATED_MOMENTS }, { "FdotAbsIntegratedMoments", GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_INTEGRATED_MOMENTS }, + { "FdotMoments", GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_MOMENTS }, + { "FdotAbsMoments", GKYL_GK_TIME_RATE_DIAGNOSTIC_FDOT_ABS_MOMENTS }, { 0, 0 } }; diff --git a/gyrokinetic/apps/gyrokinetic_multib_update_ssp_rk3.c b/gyrokinetic/apps/gyrokinetic_multib_update_ssp_rk3.c index 7c14964e6b..f3697de94e 100644 --- a/gyrokinetic/apps/gyrokinetic_multib_update_ssp_rk3.c +++ b/gyrokinetic/apps/gyrokinetic_multib_update_ssp_rk3.c @@ -379,6 +379,7 @@ gyrokinetic_multib_update_ssp_rk3(struct gkyl_gyrokinetic_multib_app* app, doubl // Compute moment of f_new to compute moment of df/dt. // Need to do it after the fields are updated. gk_species_calc_int_mom_dt(sbapp, gks, dt, gks->fdot_mom_new); + gk_species_calc_fdot_mom(sbapp, gks); } // Compute field energy divided by dt for energy balance diagnostics. diff --git a/gyrokinetic/apps/gyrokinetic_update_ssp_rk3.c b/gyrokinetic/apps/gyrokinetic_update_ssp_rk3.c index 747d9bd9bc..11ef5a782a 100644 --- a/gyrokinetic/apps/gyrokinetic_update_ssp_rk3.c +++ b/gyrokinetic/apps/gyrokinetic_update_ssp_rk3.c @@ -305,6 +305,7 @@ gyrokinetic_update_ssp_rk3(gkyl_gyrokinetic_app* app, double dt0) // Compute moment of f_new to compute moment of df/dt. // Need to do it after the fields are updated. gk_species_calc_int_mom_dt(app, gks, dt, gks->fdot_mom_new); + gk_species_calc_fdot_mom(app, gks); } // Scale species according to some criteria.