Skip to content

make the NSE solver __device__ only on GPUs #2012

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
zingale wants to merge 1 commit into
base: development
Choose a base branch
from
+12 −12
Open

make the NSE solver __device__ only on GPUs #2012

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 4 additions & 4 deletions integration/nse_update_sdc.H
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,7 +33,7 @@ using namespace amrex::literals;
/// on input, *_source are the reactive sources at time t0 and on output
/// they are the sources at time t0+dt
///
AMREX_GPU_HOST_DEVICE AMREX_INLINE
AMREX_GPU_DEVICE AMREX_INLINE
void nse_derivs(const amrex::Real rho0, const amrex::Real rhoe0, const amrex::Real *rhoaux0,
const amrex::Real dt, const amrex::Real *ydot_a,
amrex::Real& drhoedt, amrex::Real* drhoauxdt, const amrex::Real T_fixed) {
Expand Down Expand Up @@ -168,7 +168,7 @@ void nse_derivs(const amrex::Real rho0, const amrex::Real rhoe0, const amrex::Re
/// this version works with the tabulated NSE and requires AUX_THERMO
///
template <typename BurnT>
AMREX_GPU_HOST_DEVICE AMREX_INLINE
AMREX_GPU_DEVICE AMREX_INLINE
void sdc_nse_burn(BurnT& state, const amrex::Real dt) {

using namespace AuxZero;
Expand Down Expand Up @@ -307,7 +307,7 @@ void sdc_nse_burn(BurnT& state, const amrex::Real dt) {
///
/// This computes drhoedt and drhoyedt_weak for a given (rho, rho e, and rho Ye)
///
AMREX_GPU_HOST_DEVICE AMREX_INLINE
AMREX_GPU_DEVICE AMREX_INLINE
void nse_derivs(const amrex::Real rho0, const amrex::Real rhoe0,
const amrex::Real rhoYe0, amrex::Real &T0,
amrex::Real &mu_p, amrex::Real &mu_n,
Expand Down Expand Up @@ -415,7 +415,7 @@ void nse_derivs(const amrex::Real rho0, const amrex::Real rhoe0,
///

template <typename BurnT>
AMREX_GPU_HOST_DEVICE AMREX_INLINE
AMREX_GPU_DEVICE AMREX_INLINE
void sdc_nse_burn(BurnT& state, const amrex::Real dt) {

state.success = true;
Expand Down
16 changes: 8 additions & 8 deletions nse_solver/nse_check.H
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,7 @@
#include <nse_solver.H>


AMREX_GPU_HOST_DEVICE AMREX_INLINE
AMREX_GPU_DEVICE AMREX_INLINE
void check_nse_molar(const amrex::Array1D<amrex::Real, 1, NumSpec>& Y,
const amrex::Array1D<amrex::Real, 1, NumSpec>& Y_nse,
bool& nse_check) {
Expand Down Expand Up @@ -65,7 +65,7 @@ void check_nse_molar(const amrex::Array1D<amrex::Real, 1, NumSpec>& Y,
}


AMREX_GPU_HOST_DEVICE AMREX_INLINE
AMREX_GPU_DEVICE AMREX_INLINE
int get_root_index(const int nuc_idx, amrex::Array1D<int, 1, NumSpec>& group_idx) {
// Returns the root index of the nuclei given the nuclei index [1, NumSpec]
// Also updates all parent indices to point to the root index in group_idx
Expand Down Expand Up @@ -95,7 +95,7 @@ int get_root_index(const int nuc_idx, amrex::Array1D<int, 1, NumSpec>& group_idx
}


AMREX_GPU_HOST_DEVICE AMREX_INLINE
AMREX_GPU_DEVICE AMREX_INLINE
void nse_union(const int nuc_idx_a, const int nuc_idx_b,
amrex::Array1D<int, 1, NumSpec>& group_idx,
amrex::Array1D<int, 1, NumSpec>& group_size) {
Expand Down Expand Up @@ -126,7 +126,7 @@ void nse_union(const int nuc_idx_a, const int nuc_idx_b,
}


AMREX_GPU_HOST_DEVICE AMREX_INLINE
AMREX_GPU_DEVICE AMREX_INLINE
bool in_single_group(amrex::Array1D<int, 1, NumSpec>& group_idx) {

// This function checks whether all isotopes are either in the LIG group
Expand All @@ -146,7 +146,7 @@ bool in_single_group(amrex::Array1D<int, 1, NumSpec>& group_idx) {


template <typename T>
AMREX_GPU_HOST_DEVICE AMREX_INLINE
AMREX_GPU_DEVICE AMREX_INLINE
void fill_reaction_timescale(amrex::Array1D<T, 1, NSE_INDEX::NumNSERatePairs>& reaction_timescales,
const int current_rate_pair, const amrex::Real rho,
const amrex::Array1D<amrex::Real, 1, NumSpec>& Y,
Expand Down Expand Up @@ -238,7 +238,7 @@ void fill_reaction_timescale(amrex::Array1D<T, 1, NSE_INDEX::NumNSERatePairs>& r
}


AMREX_GPU_HOST_DEVICE AMREX_INLINE
AMREX_GPU_DEVICE AMREX_INLINE
void fill_merge_indices(int& merge_idx1, int& merge_idx2,
const int current_rate_pair,
amrex::Array1D<int, 1, NumSpec>& group_idx) {
Expand Down Expand Up @@ -288,7 +288,7 @@ void fill_merge_indices(int& merge_idx1, int& merge_idx2,
}


AMREX_GPU_HOST_DEVICE AMREX_INLINE
AMREX_GPU_DEVICE AMREX_INLINE
void nse_grouping(amrex::Array1D<int, 1, NumSpec>& group_idx, const amrex::Real rho,
const amrex::Array1D<amrex::Real, 1, NumSpec>& Y,
const amrex::Array1D<amrex::Real, 1, Rates::NumRates>& screened_rates,
Expand Down Expand Up @@ -374,7 +374,7 @@ void nse_grouping(amrex::Array1D<int, 1, NumSpec>& group_idx, const amrex::Real
}


AMREX_GPU_HOST_DEVICE AMREX_INLINE
AMREX_GPU_DEVICE AMREX_INLINE
bool in_nse(burn_t& current_state, bool skip_molar_check=false) {

// This function returns the boolean that tells whether we're in nse or not
Expand Down
Loading