Skip to content

Commit

Permalink
[pre-commit.ci] auto fixes from pre-commit.com hooks
Browse files Browse the repository at this point in the history
for more information, see https://pre-commit.ci
  • Loading branch information
pre-commit-ci[bot] committed Jul 22, 2024
1 parent 291b1b6 commit 94f44f3
Show file tree
Hide file tree
Showing 6 changed files with 37 additions and 41 deletions.
10 changes: 5 additions & 5 deletions src/grid/grid3D.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -8,9 +8,9 @@
#endif
#include "../global/global.h"
#include "../grid/grid3D.h"
#include "../grid/grid_enum.h" // provides grid_enum
#include "../grid/grid_enum.h" // provides grid_enum
#include "../hydro/average_cells.h" // provides Average_Slow_Cells and SlowCellConditionChecker
#include "../hydro/hydro_cuda.h" // provides Calc_dt_GPU
#include "../hydro/hydro_cuda.h" // provides Calc_dt_GPU
#include "../integrators/VL_1D_cuda.h"
#include "../integrators/VL_2D_cuda.h"
#include "../integrators/VL_3D_cuda.h"
Expand Down Expand Up @@ -449,7 +449,7 @@ void Grid3D::Execute_Hydro_Integrator(void)
#ifdef VL
VL_Algorithm_3D_CUDA(C.device, C.d_Grav_potential, H.nx, H.ny, H.nz, x_off, y_off, z_off, H.n_ghost, H.dx, H.dy,
H.dz, H.xbound, H.ybound, H.zbound, H.dt, H.n_fields, H.custom_grav, H.density_floor,
C.Grav_potential, SlowCellConditionChecker(1.0 /H.min_dt_slow,H.dx, H.dy, H.dz));
C.Grav_potential, SlowCellConditionChecker(1.0 / H.min_dt_slow, H.dx, H.dy, H.dz));
#endif // VL
#ifdef SIMPLE
Simple_Algorithm_3D_CUDA(C.device, C.d_Grav_potential, H.nx, H.ny, H.nz, x_off, y_off, z_off, H.n_ghost, H.dx, H.dy,
Expand Down Expand Up @@ -582,8 +582,8 @@ Real Grid3D::Update_Hydro_Grid()
nz_off = nz_local_start;
#endif
Average_Slow_Cells(C.device, H.nx, H.ny, H.nz, H.n_ghost, H.n_fields, gama,
SlowCellConditionChecker(max_dti_slow,H.dx, H.dy, H.dz), H.xbound,
H.ybound, H.zbound, nx_off, ny_off, nz_off);
SlowCellConditionChecker(max_dti_slow, H.dx, H.dy, H.dz), H.xbound, H.ybound, H.zbound, nx_off,
ny_off, nz_off);
#endif // AVERAGE_SLOW_CELLS

// ==Calculate the next time step using Calc_dt_GPU from hydro/hydro_cuda.h==
Expand Down
30 changes: 14 additions & 16 deletions src/hydro/average_cells.h
Original file line number Diff line number Diff line change
Expand Up @@ -10,26 +10,25 @@
#include "../global/global.h"

/*! \brief Object that checks whether a given cell meets the conditions for slow-cell averaging.
* The main motivation for creating this class is reducing ifdef statements (and allow to modify the
* actual slow-cell-condition. */
* The main motivation for creating this class is reducing ifdef statements (and allow to modify the
* actual slow-cell-condition. */
struct SlowCellConditionChecker {

// omit data-members if they aren't used for anything
#ifdef AVERAGE_SLOW_CELLS
// omit data-members if they aren't used for anything
#ifdef AVERAGE_SLOW_CELLS
Real max_dti_slow, dx, dy, dz;
#endif
#endif

/*! \brief Construct a new object. */
__host__ __device__ SlowCellConditionChecker(Real max_dti_slow, Real dx, Real dy, Real dz)
#ifdef AVERAGE_SLOW_CELLS
: max_dti_slow{max_dti_slow}, dx{dx}, dy{dy}, dz{dz}
#endif
#ifdef AVERAGE_SLOW_CELLS
: max_dti_slow{max_dti_slow}, dx{dx}, dy{dy}, dz{dz}
#endif
{
}

/*! \brief Returns whether the cell meets the condition for being considered a slow cell that must
* be averaged. */
template<bool verbose = false>
template <bool verbose = false>
__device__ bool is_slow(Real E, Real d, Real d_inv, Real vx, Real vy, Real vz, Real gamma) const
{
return this->max_dti_if_slow(E, d, d_inv, vx, vy, vz, gamma) >= 0.0;
Expand All @@ -39,18 +38,17 @@ struct SlowCellConditionChecker {
* a slow cell. If it doesn't, return a negative value instead.
*/
__device__ Real max_dti_if_slow(Real E, Real d, Real d_inv, Real vx, Real vy, Real vz, Real gamma) const;

};

#ifdef AVERAGE_SLOW_CELLS

void Average_Slow_Cells(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields,
Real gamma, SlowCellConditionChecker slow_check,
Real xbound, Real ybound, Real zbound, int nx_offset, int ny_offset, int nz_offset);
void Average_Slow_Cells(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real gamma,
SlowCellConditionChecker slow_check, Real xbound, Real ybound, Real zbound, int nx_offset,
int ny_offset, int nz_offset);

__global__ void Average_Slow_Cells_3D(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields,
Real gamma, SlowCellConditionChecker slow_check,
Real xbound, Real ybound, Real zbound, int nx_offset, int ny_offset, int nz_offset);
Real gamma, SlowCellConditionChecker slow_check, Real xbound, Real ybound,
Real zbound, int nx_offset, int ny_offset, int nz_offset);
#endif

#endif /* AVERAGE_CELLS_H */
23 changes: 11 additions & 12 deletions src/hydro/hydro_cuda.cu
Original file line number Diff line number Diff line change
Expand Up @@ -381,10 +381,11 @@ __global__ void Update_Conserved_Variables_3D(Real *dev_conserved, Real *Q_Lx, R
}
}

__global__ void PostUpdate_Conserved_Correct_Crashed_3D(Real *dev_conserved, int nx, int ny, int nz, int x_off, int y_off, int z_off,
int n_ghost, Real gamma, int n_fields, SlowCellConditionChecker slow_check)
__global__ void PostUpdate_Conserved_Correct_Crashed_3D(Real *dev_conserved, int nx, int ny, int nz, int x_off,
int y_off, int z_off, int n_ghost, Real gamma, int n_fields,
SlowCellConditionChecker slow_check)
{
int n_cells = nx * ny * nz;
int n_cells = nx * ny * nz;

// get a global thread ID
int id = threadIdx.x + blockIdx.x * blockDim.x;
Expand All @@ -394,9 +395,8 @@ __global__ void PostUpdate_Conserved_Correct_Crashed_3D(Real *dev_conserved, int

if (xid > n_ghost - 1 && xid < nx - n_ghost && yid > n_ghost - 1 && yid < ny - n_ghost && zid > n_ghost - 1 &&
zid < nz - n_ghost) {

#if !(defined(DENSITY_FLOOR) && defined(TEMPERATURE_FLOOR))
// threads corresponding to real cells do the calculation
// threads corresponding to real cells do the calculation
if (dev_conserved[id] < 0.0 || dev_conserved[id] != dev_conserved[id] || dev_conserved[4 * n_cells + id] < 0.0 ||
dev_conserved[4 * n_cells + id] != dev_conserved[4 * n_cells + id]) {
printf("%3d %3d %3d Thread crashed in final update. %e - - - %e\n", xid + x_off, yid + y_off, zid + z_off,
Expand Down Expand Up @@ -683,7 +683,8 @@ void Temperature_Ceiling(Real *dev_conserved, int nx, int ny, int nz, int n_ghos
}
}

__device__ Real SlowCellConditionChecker::max_dti_if_slow(Real E, Real d, Real d_inv, Real vx, Real vy, Real vz, Real gamma) const
__device__ Real SlowCellConditionChecker::max_dti_if_slow(Real E, Real d, Real d_inv, Real vx, Real vy, Real vz,
Real gamma) const
{
#ifndef AVERAGE_SLOW_CELLS
return -1.0;
Expand All @@ -695,9 +696,9 @@ __device__ Real SlowCellConditionChecker::max_dti_if_slow(Real E, Real d, Real d

#ifdef AVERAGE_SLOW_CELLS

void Average_Slow_Cells(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields,
Real gamma, SlowCellConditionChecker slow_check,
Real xbound, Real ybound, Real zbound, int nx_offset, int ny_offset, int nz_offset)
void Average_Slow_Cells(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real gamma,
SlowCellConditionChecker slow_check, Real xbound, Real ybound, Real zbound, int nx_offset,
int ny_offset, int nz_offset)
{
// set values for GPU kernels
int n_cells = nx * ny * nz;
Expand Down Expand Up @@ -1309,9 +1310,7 @@ __device__ void Average_Cell_All_Fields(int i, int j, int k, int nx, int ny, int
for (int kk = k - 1; kk <= k + 1; kk++) {
for (int jj = j - 1; jj <= j + 1; jj++) {
for (int ii = i - 1; ii <= i + 1; ii++) {

if (ii <= stale_depth - 1 || ii >= nx - stale_depth ||
jj <= stale_depth - 1 || jj >= ny - stale_depth ||
if (ii <= stale_depth - 1 || ii >= nx - stale_depth || jj <= stale_depth - 1 || jj >= ny - stale_depth ||
kk <= stale_depth - 1 || kk >= nz - stale_depth) {
continue;
}
Expand Down
7 changes: 3 additions & 4 deletions src/hydro/hydro_cuda.h
Original file line number Diff line number Diff line change
Expand Up @@ -22,8 +22,9 @@ __global__ void Update_Conserved_Variables_3D(Real *dev_conserved, Real *Q_Lx, R
Real gamma, int n_fields, int custom_grav, Real density_floor,
Real *dev_potential);

__global__ void PostUpdate_Conserved_Correct_Crashed_3D(Real *dev_conserved, int nx, int ny, int nz, int x_off, int y_off, int z_off,
int n_ghost, Real gamma, int n_fields, SlowCellConditionChecker slow_check);
__global__ void PostUpdate_Conserved_Correct_Crashed_3D(Real *dev_conserved, int nx, int ny, int nz, int x_off,
int y_off, int z_off, int n_ghost, Real gamma, int n_fields,
SlowCellConditionChecker slow_check);

/*!
* \brief Determine the maximum inverse crossing time in a specific cell
Expand Down Expand Up @@ -84,8 +85,6 @@ void Temperature_Ceiling(Real *dev_conserved, int nx, int ny, int nz, int n_ghos
Real T_ceiling);
#endif // TEMPERATURE CEILING



void Apply_Temperature_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real U_floor);

__global__ void Temperature_Floor_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields,
Expand Down
6 changes: 3 additions & 3 deletions src/integrators/VL_3D_cuda.cu
Original file line number Diff line number Diff line change
Expand Up @@ -36,7 +36,7 @@ __global__ void Update_Conserved_Variables_3D_half(Real *dev_conserved, Real *de
void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int ny, int nz, int x_off, int y_off,
int z_off, int n_ghost, Real dx, Real dy, Real dz, Real xbound, Real ybound, Real zbound,
Real dt, int n_fields, int custom_grav, Real density_floor, Real *host_grav_potential,
const SlowCellConditionChecker& slow_check)
const SlowCellConditionChecker &slow_check)
{
// Here, *dev_conserved contains the entire
// set of conserved variables on the grid
Expand Down Expand Up @@ -376,8 +376,8 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int

// Step 6b: Address any crashed threads
hipLaunchKernelGGL(PostUpdate_Conserved_Correct_Crashed_3D, update_full_launch_params.get_numBlocks(),
update_full_launch_params.get_threadsPerBlock(), 0, 0, dev_conserved, nx, ny, nz, x_off, y_off, z_off,
n_ghost, gama, n_fields, slow_check);
update_full_launch_params.get_threadsPerBlock(), 0, 0, dev_conserved, nx, ny, nz, x_off, y_off,
z_off, n_ghost, gama, n_fields, slow_check);

#ifdef MHD
// Update the magnetic fields
Expand Down
2 changes: 1 addition & 1 deletion src/integrators/VL_3D_cuda.h
Original file line number Diff line number Diff line change
Expand Up @@ -10,7 +10,7 @@
void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int ny, int nz, int x_off, int y_off,
int z_off, int n_ghost, Real dx, Real dy, Real dz, Real xbound, Real ybound, Real zbound,
Real dt, int n_fields, int custom_grav, Real density_floor, Real *host_grav_potential,
const SlowCellConditionChecker& slow_check);
const SlowCellConditionChecker &slow_check);

void Free_Memory_VL_3D();

Expand Down

0 comments on commit 94f44f3

Please sign in to comment.