Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
109 commits
Select commit Hold shift + click to select a range
f314265
Initial commit. Kernels and updaters for calculating drag and diffusi…
May 4, 2023
f86a254
Cleanup of unit test
May 4, 2023
f11f35b
Infrastructure to project Maxwellian H and G potentials onto basis fu…
May 10, 2023
e2d9cc6
Quotes to <> in some includes and added release for projection object
May 10, 2023
5f1881f
Updater to project analytic Rosenbluth potentials for a Maxwellian. P…
Jul 6, 2023
657b12e
Current state of proj_maxwellian_pots_on_basis: quadrature projection…
Nov 3, 2023
3d06622
Projecting potentials onto nodes rather than with quadrature. Continu…
Nov 4, 2023
3f8e05d
Now when projecting Maxwellian potentials, use proj_on_basis routine …
Jan 4, 2024
3839212
Drag and diffusion coefficient calculation now working and converging…
Jan 4, 2024
3ae14b0
Kernels and updaters for FPO drag and diffusion contributions. Also u…
Feb 16, 2024
34a5e0e
Drag term update uses gen_stencil_advance like the diffusion term.
Feb 16, 2024
3290833
Skeleton of FPO app
JunoRavin Feb 16, 2024
30e6c32
Adding identical tests to LBO relaxation in 1x3v, but for FPO
JunoRavin Feb 16, 2024
81135c5
Merge branch 'main' into fpo-john
JunoRavin Feb 16, 2024
551469d
Many updates to FPO, still testing to be sure but it seems to be work…
Mar 11, 2024
56c128c
Renaming potential projection files. Unit test for these is incoming
Mar 11, 2024
fe64b61
Unit test for potential projections, plus it works for hybrid basis n…
Mar 12, 2024
d6250a1
Calculate prim vars in FPO update with gkyl_dg_prim_vars updater rath…
Mar 16, 2024
dd3551d
Revert "Calculate prim vars in FPO update with gkyl_dg_prim_vars upda…
Mar 16, 2024
a278bdb
Add function to project analytic dH/dv and d2G/dv2 for cross species …
Mar 20, 2024
3687d15
Add checks for derivative values to the unit test
Mar 20, 2024
97c05b8
Merge branch 'main' into fpo-john
JunoRavin Apr 4, 2024
7cb259a
Commiting a first pass at the needed moments in FPO to create a conse…
JunoRavin Apr 5, 2024
1d5c3a1
Getting rid of some variable-size initialization that is not needed
JunoRavin Apr 5, 2024
6ba0446
Need to initialize quantities in the hyper_dg_gen_stencil, otherwise …
JunoRavin Apr 5, 2024
ebd98e0
Small clean up of the gen_stencil object to not use variables we don'…
JunoRavin Apr 5, 2024
13eb0fc
getting rid of the initialization of these now unneeded variables
JunoRavin Apr 5, 2024
f56feaa
Some clean up in preparation for GPU-ification. Removing the lab_moms…
JunoRavin Apr 6, 2024
298a4ae
Some more changes to the regression tests to facilitate easier testin…
JunoRavin Apr 6, 2024
44b7e43
Adding the ground work for GPU-ification of FPO solver. We will still…
JunoRavin Apr 6, 2024
61cc3f0
Fixing some GPU-specific bugs. Also required fixing a warning in the …
JunoRavin Apr 6, 2024
21fe29f
Changing some names in the FPO drag to be more consistent with what o…
JunoRavin Apr 6, 2024
1298014
Range is passed by value to CUDA kernels so need to be careful about …
JunoRavin Apr 6, 2024
d3f1295
Need to compile the FPO kernels with the -x flag so they can be linke…
JunoRavin Apr 6, 2024
1883b67
Removing an old file before I forget
Apr 11, 2024
be77e24
Value for velocity at domain boundary for FPO potential surface proje…
Apr 14, 2024
d29e92f
Changed the drag term for FPO to use the same technique as the canoni…
May 6, 2024
0d34158
Updated how stenciling works for hyper_dg_gen_stencil and the diffusi…
May 23, 2024
9fbc300
Diffusion tensor calculation now uses the same indexing as the diffus…
May 27, 2024
98a6c11
Added calls to kernels to calculate boundary corrections for FPO mome…
May 31, 2024
b8b48ad
Hooked in calculation of correction volume and boundary moments, line…
Jun 27, 2024
2c2c2a9
Merge branch 'main' into fpo-john
Jul 2, 2024
e838570
Updates to boundary corrections for momentum conservation constraint …
Jul 17, 2024
e6a1fc4
Whoops forgot some files
Jul 17, 2024
09c1fd2
Cleaning up old debug code that broke something
Jul 17, 2024
cfe374b
Merge branch 'main' into fpo-john
Aug 28, 2024
5ed2be8
First pass at finishing up the GPU support for FPO. Probably issues b…
Sep 11, 2024
447733a
Continuing to add GPU support to FPO. Fixed a bug in potential projec…
Sep 16, 2024
e25774c
Forgot to swap __device__ for gkyl macro in C header file
Sep 16, 2024
a2bdcef
Regenerated kernels so stencil indexing is correct again
Sep 16, 2024
3df6f2e
Missed header file
Sep 16, 2024
8859c89
Fixed some bugs and formatting. Added CUDA kernels for coefficient co…
Sep 19, 2024
30397bb
Fixed an indexing issue, full FPO runs well on CPUs. Added those rege…
Sep 19, 2024
5cf7801
This should have been a few commits but things kept breaking. Added G…
Oct 4, 2024
5bd8c5b
Fix potentials release function bug for CPUs
Oct 4, 2024
909c85d
Found an issue I had introdued in the diffusion kernels that broke co…
Oct 7, 2024
caf01a0
oops broke GPUs again before, this potential projection object needs …
Oct 7, 2024
81e7a66
Fixed some inconsistencies in boundary correction calculation, now FP…
Oct 13, 2024
69befa1
Oops missed a change
Oct 13, 2024
701e2a5
Merge branch 'main' into fpo-john
Dec 18, 2024
0e2dca3
Unit test for FPO calculation of potentials and drag/diffusion coeffi…
Dec 27, 2024
152d940
Forgot to be good and release memory. Also GPU test
Dec 28, 2024
309d9a7
Updating some kernels. Some p1 kernels are old (because testing was a…
Jan 3, 2025
aab88cf
Some cleanup and p1 hybrid basis support. Does not conserve energy fo…
Jan 24, 2025
c9c3fe7
Renaming conservation correction routines because it has bugged me
Jan 24, 2025
1edaec9
Update unit test to reflect changes to diffusion coefficient kernels
Jan 29, 2025
9389c66
Unit test for checking conservation correction moments
Jan 30, 2025
55d3879
Integrated moment kernels used pure serendipity basis for p1 rather t…
Jan 30, 2025
0ea670c
Added support for 2x3v p1 FPO
Feb 2, 2025
97dcfe9
Changes to kernel selection that hopefully fixes const initialization…
Feb 3, 2025
90c7085
Adding GPU support for p1 1x3v and 2x3v
Feb 5, 2025
66a071f
Sneaky bug in conservation correction routine. Now 2x3v actually work…
Feb 6, 2025
d2fbf2f
Renamed a factor in the maxima scripts for the drag/diffusion terms a…
Feb 8, 2025
1850250
Precision change to unit test
Feb 8, 2025
ec729d9
Commenting out the write function to output H, G, a and D
Feb 8, 2025
314db40
Instead, write method for FPO quantities can now be requested by addi…
Feb 8, 2025
c5b314b
Merge branch 'main' into fpo-john. Just a few commits to bring into t…
JunoRavin Feb 12, 2025
9b159e4
Refactoring FPO regression test to follow new regression test syntax …
JunoRavin Feb 12, 2025
4cb8532
Separating boundary correction calculation kernels so they aren't so …
Feb 13, 2025
4f904d4
Nevermind, instead changing bcorr kernels slightly to shorten line le…
Feb 13, 2025
afe196d
Missed 2x3v
Feb 13, 2025
15df83f
Missed a factoring constant
Feb 13, 2025
696414c
Some restructuring of conservation moment kernels to reduce line leng…
Feb 13, 2025
ed35faf
The FPO branch is encountering a truly agonizingly annoying GPU bug w…
JunoRavin Feb 14, 2025
5679ca2
Merge branch 'main' into fpo-john
JunoRavin Apr 10, 2025
892c436
Merge branch 'main' into fpo-john
JunoRavin May 6, 2025
5a1ca83
Fixing the hyper dg regression test
JunoRavin May 6, 2025
d5675da
Bug introduced some time ago during a merge from main, a member of th…
jmrodman Jun 26, 2025
966e396
Merge branch 'main' into fpo-john. Somewhat non-trivial due to how gi…
JunoRavin Jul 9, 2025
f3ba9ff
In anticipation of the fact that we are not going to support p>1 in c…
JunoRavin Jul 9, 2025
60b8e2f
Merge branch 'main' into fpo-john
JunoRavin Oct 20, 2025
49d2e8b
Update to Rosenbluth potentials form that somehow avoids singularity …
jmrodman Jan 12, 2026
721114a
Regenerated all the FPO kernels. Swapped p=1 from hybrid back to pure…
jmrodman Jan 25, 2026
511b6ae
Merge branch 'main' into fpo-john
jmrodman Jan 25, 2026
1fd6cdf
Missed a kernel. Fixed some header files so we correctly load p2 kernels
jmrodman Jan 26, 2026
1a49c30
Cleaning up old diffusion tensor files. Naming convention changed sli…
jmrodman Jan 26, 2026
625171b
Fixed an error in the Maxima factoring script that made for some empt…
jmrodman Jan 26, 2026
0e26494
Bug in regression test: one distribution was using LBO
jmrodman Jan 26, 2026
8966234
Oops p=1 was still using hybrid because of Maxima defaulting to it wi…
jmrodman Jan 28, 2026
5951e92
Fixing dumb indexing error in upper off-diagonal terms
jmrodman Jan 28, 2026
9ad4d24
Apparently the inconsistency between indexing on CPUs and GPUs in mom…
jmrodman Jan 29, 2026
ad6a61a
Committed the wrong kernels for the diffusion surface update
jmrodman Jan 30, 2026
651c8c5
aaaaaand the remaining diffusion surface term kernels
jmrodman Jan 30, 2026
e9ef0ad
Adjusting unit test values and structure for drag/diffusion coefficie…
jmrodman Jan 30, 2026
dcf0caf
Using maxima-generated header files now
jmrodman Jan 30, 2026
be8287b
Merge branch 'main' into fpo-john
jmrodman Feb 2, 2026
bbd514d
Merge branch 'main' into fpo-john
jmrodman Mar 12, 2026
a7148b3
Merge branch 'main' into fpo-john
jmrodman Mar 16, 2026
4ca1294
Merge branch 'main' into fpo-john
jmrodman Apr 22, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
The table of contents is too big for display.
Diff view
Diff view
  •  
  •  
  •  
1,919 changes: 0 additions & 1,919 deletions gyrokinetic/ker/lbo_gyrokinetic/gyrokinetic_self_prim_moments_3x2v_ser_p2.c

This file was deleted.

4 changes: 2 additions & 2 deletions gyrokinetic/zero/gkyl_prim_lbo_gyrokinetic_priv.h
Original file line number Diff line number Diff line change
Expand Up @@ -31,7 +31,7 @@ static const gkyl_prim_lbo_gyrokinetic_kern_list ser_self_prim_kernels[] = {
// 2x kernels
{ NULL, gyrokinetic_self_prim_moments_2x2v_ser_p1, gyrokinetic_self_prim_moments_2x2v_ser_p2 }, // 2
// 3x kernels
{ NULL, gyrokinetic_self_prim_moments_3x2v_ser_p1, gyrokinetic_self_prim_moments_3x2v_ser_p2 }, // 3
{ NULL, gyrokinetic_self_prim_moments_3x2v_ser_p1, NULL }, // 3
};

// cross primitive moment kernel list
Expand All @@ -43,7 +43,7 @@ static const gkyl_prim_lbo_gyrokinetic_cross_kern_list ser_cross_prim_kernels[]
// 2x kernels
{ NULL, gyrokinetic_cross_prim_moments_2x2v_ser_p1, gyrokinetic_cross_prim_moments_2x2v_ser_p2 }, // 2
// 3x kernels
{ NULL, gyrokinetic_cross_prim_moments_3x2v_ser_p1, gyrokinetic_cross_prim_moments_3x2v_ser_p2 }, // 3
{ NULL, gyrokinetic_cross_prim_moments_3x2v_ser_p1, NULL }, // 3
};

struct prim_lbo_type_gyrokinetic {
Expand Down
11 changes: 11 additions & 0 deletions vlasov/apps/gkyl_vlasov.h
Original file line number Diff line number Diff line change
Expand Up @@ -45,6 +45,7 @@ struct gkyl_vlasov_projection {
// Parameters for species collisions
struct gkyl_vlasov_collisions {
enum gkyl_collision_id collision_id; // type of collisions (see gkyl_eqn_type.h)
bool write_diagnostics; // Whether to write diagnostics out.

void *ctx; // context for collision function
// function for computing self-collision frequency
Expand Down Expand Up @@ -552,6 +553,16 @@ void gkyl_vlasov_app_write_species(gkyl_vlasov_app* app, int sidx, double tm, in
*/
void gkyl_vlasov_app_write_species_lte(gkyl_vlasov_app* app, int sidx, double tm, int frame);

/**
* Write species data to file for H, G, drag coeff, diff coeff.
*
* @param app App object.
* @param sidx Index of species to initialize.
* @param tm Time-stamp
* @param frame Frame number
*/
void gkyl_vlasov_app_write_species_fpo(gkyl_vlasov_app* app, int sidx, double tm, int frame);

/**
* Write fluid species data to file.
*
Expand Down
90 changes: 90 additions & 0 deletions vlasov/apps/gkyl_vlasov_priv.h
Original file line number Diff line number Diff line change
Expand Up @@ -34,11 +34,14 @@
#include <gkyl_dg_canonical_pb.h>
#include <gkyl_dg_canonical_pb_fluid.h>
#include <gkyl_dg_euler.h>
#include <gkyl_dg_fpo_vlasov_diff_coeff.h>
#include <gkyl_dg_fpo_vlasov_drag_coeff.h>
#include <gkyl_dg_maxwell.h>
#include <gkyl_dg_updater_fluid.h>
#include <gkyl_dg_updater_diffusion_fluid.h>
#include <gkyl_dg_updater_diffusion_gen.h>
#include <gkyl_dg_updater_lbo_vlasov.h>
#include <gkyl_dg_updater_fpo_vlasov.h>
#include <gkyl_dg_updater_rad_vlasov.h>
#include <gkyl_dg_updater_moment.h>
#include <gkyl_dg_updater_vlasov.h>
Expand All @@ -49,13 +52,17 @@
#include <gkyl_dynvec.h>
#include <gkyl_elem_type.h>
#include <gkyl_eqn_type.h>
#include <gkyl_fpo_vlasov_coeff_correct.h>
#include <gkyl_fpo_vlasov_coeff_recovery.h>
#include <gkyl_fpo_proj_maxwellian_pots_on_basis.h>
#include <gkyl_eval_on_nodes.h>
#include <gkyl_fem_poisson.h>
#include <gkyl_ghost_surf_calc.h>
#include <gkyl_hyper_dg.h>
#include <gkyl_mom_bcorr_lbo_vlasov.h>
#include <gkyl_mom_calc.h>
#include <gkyl_mom_calc_bcorr.h>
#include <gkyl_mom_fpo_vlasov.h>
#include <gkyl_mom_vlasov.h>
#include <gkyl_mom_vlasov_sr.h>
#include <gkyl_null_pool.h>
Expand Down Expand Up @@ -205,6 +212,40 @@ struct vm_rad_drag {
gkyl_dg_updater_rad_vlasov *rad_slvr; // radiation solver
};

struct vm_fpo_collisions {
bool write_diagnostics; // Whether to write diagnostics out.
struct gkyl_array *gamma; // FPO Gamma factor
struct gkyl_array *h, *g; // Rosenbluth potentials
struct gkyl_array *h_host, *g_host;

// Maxwellian potentials and derivatives on velocity space edges for boundary conditions
struct gkyl_array *h_surf, *g_surf;
struct gkyl_array *dhdv_surf, *dgdv_surf, *d2gdv2_surf; // dG/dv evaluated at transverse surfaces, dH/dv and d2G/dv2 at same surface

gkyl_proj_maxwellian_pots_on_basis *pot_slvr; // potential solver for Maxwellian potentials

struct vm_species_moment lte_moms; // calculator for LTE moments for potential calculation (n u_drift, T/m)
struct vm_species_moment moms; // calculator for moments (M0, M1i, M2)

struct gkyl_array *drag_coeff, *diff_coeff; // Drag and diffusion coefficients
struct gkyl_array *drag_coeff_host, *diff_coeff_host; // Drag and diffusion coefficients
struct gkyl_array *drag_coeff_surf, *diff_coeff_surf; // Drag and diffusion coefficient surface expansions at lower cell boundary
struct gkyl_array *sgn_drag_coeff_surf, *const_sgn_drag_coeff_surf; // Sign of drag coefficient at lower boundary of cell, and boolean for checking if sign(drag_coeff) is constant along boundary

gkyl_fpo_vlasov_coeff_recovery *coeff_recovery; // Struct for drag and diffusion coeff calculation

// Momentum and energy conservation corrections
struct gkyl_array *fpo_moms, *boundary_corrections; // Primitive moments and boundary corrections
struct gkyl_array *drag_diff_coeff_corrs; // Correction quantities added to drag and diffusion coefficients

struct gkyl_mom_calc *fpo_mom_calc; // FPO volume corrections calculator
struct gkyl_mom_calc_bcorr *bcorr_calc; // FPO boundary corrections calculator
gkyl_fpo_coeff_correct *coeff_correct_calc; // FPO drag and diffusion coeff correction calculator

gkyl_dg_updater_collisions *coll_slvr; // collision solver
long offsets[36]; // Array of relative offsets for 3- and 6-cell stencils. Will probably be moved to hyper_dg
};

struct vm_boundary_fluxes {
struct gkyl_rect_grid boundary_grid[2*GKYL_MAX_CDIM];
struct gkyl_array *flux_arr[2*GKYL_MAX_CDIM];
Expand Down Expand Up @@ -429,6 +470,9 @@ struct vm_species {
struct {
struct vm_bgk_collisions bgk; // BGK collisions object
};
struct {
struct vm_fpo_collisions fpo; // FPO collisions object
};
};

enum gkyl_radiation_id radiation_id; // type of radiation
Expand Down Expand Up @@ -1090,6 +1134,52 @@ void vm_species_bgk_rhs(gkyl_vlasov_app *app,
*/
void vm_species_bgk_release(const struct gkyl_vlasov_app *app, const struct vm_bgk_collisions *bgk);

/** vm_species_fpo API */

/**
* Initialize species fpo collisions object.
*
* @param app Vlasov app object
* @param s Species object
* @param fpo Species fpo object
* @param collides_with_fluid Boolean for if kinetic species collides with a fluid species
*/
void vm_species_fpo_init(struct gkyl_vlasov_app *app, struct vm_species *s,
struct vm_fpo_collisions *fpo);

/**
* Compute drag and diffusion coefficients for update
*
* @param app Vlasov app object
* @param species Pointer to species
* @param fpo Pointer to fpo
* @param fin Input distribution function
*/
void vm_species_fpo_drag_diff_coeffs(gkyl_vlasov_app *app, const struct vm_species *s,
struct vm_fpo_collisions *fpo, const struct gkyl_array *fin);

/**
* Compute RHS from fpo collisions
*
* @param app Vlasov app object
* @param species Pointer to species
* @param fpo Pointer to fpo
* @param fin Input distribution function
* @param rhs On output, the RHS from fpo
* @return Maximum stable time-step
*/
void vm_species_fpo_rhs(gkyl_vlasov_app *app,
const struct vm_species *species,
struct vm_fpo_collisions *fpo,
const struct gkyl_array *fin, struct gkyl_array *rhs);

/**
* Release species fpo object.
*
* @param app Vlasov app object
* @param sm Species fpo object to release
*/
void vm_species_fpo_release(const struct gkyl_vlasov_app *app, const struct vm_fpo_collisions *fpo);
/** vm_species_radiation API */

/**
Expand Down
56 changes: 56 additions & 0 deletions vlasov/apps/vlasov.c
Original file line number Diff line number Diff line change
Expand Up @@ -610,6 +610,9 @@ gkyl_vlasov_app_write(gkyl_vlasov_app* app, double tm, int frame)
if (app->species[i].info.output_f_lte) {
gkyl_vlasov_app_write_species_lte(app, i, tm, frame);
}
if ((app->species[i].collision_id == GKYL_FPO_COLLISIONS) && (app->species[i].fpo.write_diagnostics)) {
gkyl_vlasov_app_write_species_fpo(app, i, tm, frame);
}
}
for (int i=0; i<app->num_fluid_species; ++i) {
gkyl_vlasov_app_write_fluid_species(app, i, tm, frame);
Expand Down Expand Up @@ -781,6 +784,59 @@ gkyl_vlasov_app_write_species_lte(gkyl_vlasov_app* app, int sidx, double tm, int
vlasov_array_meta_release(mt);
}

void
gkyl_vlasov_app_write_species_fpo(gkyl_vlasov_app* app, int sidx, double tm, int frame)
{
struct gkyl_msgpack_data *mt = vlasov_array_meta_new( (struct vlasov_output_meta) {
.frame = frame,
.stime = tm,
.poly_order = app->poly_order,
.basis_type = app->basis.id
}
);

struct vm_species *vm_s = &app->species[sidx];

const char *fmt_h = "%s-%s_H_%d.gkyl";
const char *fmt_g = "%s-%s_G_%d.gkyl";
const char *fmt_drag = "%s-%s_drag_coeff_%d.gkyl";
const char *fmt_diff = "%s-%s_diff_coeff_%d.gkyl";
int sz = gkyl_calc_strlen(fmt_drag, app->name,vm_s->info.name, frame);
char fileNm[sz+1]; // ensures no buffer overflow

vm_species_fpo_drag_diff_coeffs(app, vm_s, &vm_s->fpo, vm_s->f);

if (app->use_gpu) {
// copy data from device to host before writing it out
gkyl_array_copy(vm_s->fpo.h_host, vm_s->fpo.h);
gkyl_array_copy(vm_s->fpo.g_host, vm_s->fpo.g);
gkyl_array_copy(vm_s->fpo.drag_coeff_host, vm_s->fpo.drag_coeff);
gkyl_array_copy(vm_s->fpo.diff_coeff_host, vm_s->fpo.diff_coeff);
}

// Write H
snprintf(fileNm, sizeof fileNm, fmt_h, app->name,vm_s->info.name, frame);
gkyl_comm_array_write(vm_s->comm, &vm_s->grid, &vm_s->local,
mt, vm_s->fpo.h_host, fileNm);

// Write G
snprintf(fileNm, sizeof fileNm, fmt_g, app->name,vm_s->info.name, frame);
gkyl_comm_array_write(vm_s->comm, &vm_s->grid, &vm_s->local,
mt, vm_s->fpo.g_host, fileNm);

// Write drag coefficient
snprintf(fileNm, sizeof fileNm, fmt_drag, app->name,vm_s->info.name, frame);
gkyl_comm_array_write(vm_s->comm, &vm_s->grid, &vm_s->local,
mt, vm_s->fpo.drag_coeff_host, fileNm);

// Write diffusion tensor
snprintf(fileNm, sizeof fileNm, fmt_diff, app->name,vm_s->info.name, frame);
gkyl_comm_array_write(vm_s->comm, &vm_s->grid, &vm_s->local,
mt, vm_s->fpo.diff_coeff_host, fileNm);

vlasov_array_meta_release(mt);
}

void
gkyl_vlasov_app_write_fluid_species(gkyl_vlasov_app* app, int sidx, double tm, int frame)
{
Expand Down
4 changes: 4 additions & 0 deletions vlasov/apps/vlasov_forward_euler.c
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,10 @@ vlasov_forward_euler(gkyl_vlasov_app* app, double tcurr, double dt,
vm_species_bgk_moms(app, &app->species[i],
&app->species[i].bgk, fin[i]);
}
else if (app->species[i].collision_id == GKYL_FPO_COLLISIONS) {
vm_species_fpo_drag_diff_coeffs(app, &app->species[i],
&app->species[i].fpo, fin[i]);
}
}

// compute necessary moments for cross-species collisions
Expand Down
14 changes: 12 additions & 2 deletions vlasov/apps/vm_species.c
Original file line number Diff line number Diff line change
Expand Up @@ -225,7 +225,7 @@ vm_species_init(struct gkyl_vm *vm, struct gkyl_vlasov_app *app, struct vm_speci
s->sgn_alpha_surf = mkarr(app->use_gpu, sgn_alpha_surf_sz, s->local_ext.volume);
s->const_sgn_alpha = mk_int_arr(app->use_gpu, (cdim + vdim), s->local_ext.volume);

// Pre-compute alpha_surf, sgn_alpha_surf, const_sgn_alpha, and cot_vec since they are time-independent
// Pre-compute alpha_surf, sgn_alpha_surf, and const_sgn_alpha since they are time-independent
struct gkyl_dg_calc_canonical_pb_vars *calc_vars = gkyl_dg_calc_canonical_pb_vars_new(&s->grid,
&app->confBasis, &app->basis, app->use_gpu);
gkyl_dg_calc_canonical_pb_vars_alpha_surf(calc_vars, &app->local, &s->local, &s->local_ext, s->hamil,
Expand All @@ -241,7 +241,7 @@ vm_species_init(struct gkyl_vm *vm, struct gkyl_vlasov_app *app, struct vm_speci
}
else {
if (s->field_id == GKYL_FIELD_NULL || s->field_id == GKYL_FIELD_E_B) {
struct gkyl_dg_vlasov_auxfields aux_inp = {.field = s->qmem, .cot_vec = 0,
struct gkyl_dg_vlasov_auxfields aux_inp = {.field = s->qmem,
.alpha_surf = 0, .sgn_alpha_surf = 0, .const_sgn_alpha = 0 };
s->slvr = gkyl_dg_updater_vlasov_new(&s->grid, &app->confBasis, &app->basis,
&app->local, &s->local_vel, &s->local, is_zero_flux, s->model_id, s->field_id, &aux_inp, app->use_gpu);
Expand Down Expand Up @@ -327,6 +327,7 @@ vm_species_init(struct gkyl_vm *vm, struct gkyl_vlasov_app *app, struct vm_speci
// initialize empty collision structs so inputs of structs are set to 0
s->lbo = (struct vm_lbo_collisions) { };
s->bgk = (struct vm_bgk_collisions) { };
s->fpo = (struct vm_fpo_collisions) { };
if (s->info.output_f_lte){
// Always have correct moments on for the f_lte output
struct correct_all_moms_inp corr_inp = { .correct_all_moms = true,
Expand All @@ -340,6 +341,9 @@ vm_species_init(struct gkyl_vm *vm, struct gkyl_vlasov_app *app, struct vm_speci
else if (s->collision_id == GKYL_BGK_COLLISIONS) {
vm_species_bgk_init(app, s, &s->bgk);
}
else if (s->collision_id == GKYL_FPO_COLLISIONS) {
vm_species_fpo_init(app, s, &s->fpo);
}

// determine radiation type to use in vlasov update
s->radiation_id = s->info.radiation.radiation_id;
Expand Down Expand Up @@ -515,6 +519,9 @@ vm_species_rhs(gkyl_vlasov_app *app, struct vm_species *species,
species->bgk.implicit_step = false;
vm_species_bgk_rhs(app, species, &species->bgk, fin, rhs);
}
else if (species->collision_id == GKYL_FPO_COLLISIONS) {
vm_species_fpo_rhs(app, species, &species->fpo, fin, rhs);
}

if (species->calc_bflux) {
vm_species_bflux_rhs(app, species, &species->bflux, fin, rhs);
Expand Down Expand Up @@ -810,6 +817,9 @@ vm_species_release(const gkyl_vlasov_app* app, const struct vm_species *s)
else if (s->collision_id == GKYL_BGK_COLLISIONS) {
vm_species_bgk_release(app, &s->bgk);
}
else if (s->collision_id == GKYL_FPO_COLLISIONS) {
vm_species_fpo_release(app, &s->fpo);
}

if (s->radiation_id == GKYL_VM_COMPTON_RADIATION) {
vm_species_radiation_release(app, &s->rad);
Expand Down
Loading
Loading