Skip to content

Commit

Permalink
Merge pull request #203 from alwinm/main-float32
Browse files Browse the repository at this point in the history
Float 32 output, HDF5 overhaul, static_grav and other minor fixes. (copy for main)
  • Loading branch information
evaneschneider authored Oct 25, 2022
2 parents cf898d5 + 03b3482 commit d2a851d
Show file tree
Hide file tree
Showing 12 changed files with 578 additions and 570 deletions.
31 changes: 31 additions & 0 deletions builds/make.type.rot_proj
Original file line number Diff line number Diff line change
@@ -0,0 +1,31 @@
#-- Default hydro only build with rotated projection

DFLAGS += -DCUDA
DFLAGS += -DMPI_CHOLLA
DFLAGS += -DPRECISION=2
DFLAGS += -DPPMC
DFLAGS += -DHLLC

# Integrator
DFLAGS += -DSIMPLE
#DFLAGS += -DVL

# Apply a density and temperature floor
DFLAGS += -DDENSITY_FLOOR
DFLAGS += -DTEMPERATURE_FLOOR

# Solve the Gas Internal Energy usisng a Dual Energy Formalism
#DFLAGS += -DDE

# Apply cooling on the GPU from precomputed tables
#DFLAGS += -DCOOLING_GPU

# Measure the Timing of the different stages
#DFLAGS += -DCPU_TIME

# Select output format
# Can also add -DSLICES and -DPROJECTIONS
OUTPUT ?= -DOUTPUT -DHDF5
DFLAGS += $(OUTPUT)

DFLAGS += -DROTATED_PROJECTION
32 changes: 32 additions & 0 deletions builds/make.type.static_grav
Original file line number Diff line number Diff line change
@@ -0,0 +1,32 @@
#-- Default hydro only build with static_grav

DFLAGS += -DCUDA
DFLAGS += -DMPI_CHOLLA
DFLAGS += -DPRECISION=2
DFLAGS += -DPPMC
DFLAGS += -DHLLC

# Integrator
DFLAGS += -DSIMPLE
#DFLAGS += -DVL

# Apply a density and temperature floor
DFLAGS += -DDENSITY_FLOOR
DFLAGS += -DTEMPERATURE_FLOOR

# Solve the Gas Internal Energy usisng a Dual Energy Formalism
#DFLAGS += -DDE

DFLAGS += -DSTATIC_GRAV

# Apply cooling on the GPU from precomputed tables
#DFLAGS += -DCOOLING_GPU

# Measure the Timing of the different stages
#DFLAGS += -DCPU_TIME

# Select output format
# Can also add -DSLICES and -DPROJECTIONS
OUTPUT ?= -DOUTPUT -DHDF5
DFLAGS += $(OUTPUT)

59 changes: 59 additions & 0 deletions examples/3D/float32_sound_wave.txt
Original file line number Diff line number Diff line change
@@ -0,0 +1,59 @@
#
# Parameter File for sound wave test with float32 output
#

################################################
# number of grid cells in the x dimension
nx=256
# number of grid cells in the y dimension
ny=256
# number of grid cells in the z dimension
nz=256
# final output time
tout=0.05
# time interval for output
outstep=0.05
# name of initial conditions
init=Sound_Wave
# domain properties
xmin=0.0
ymin=0.0
zmin=0.0
xlen=4.0
ylen=4.0
zlen=4.0
# type of boundary conditions
xl_bcnd=1
xu_bcnd=1
yl_bcnd=1
yu_bcnd=1
zl_bcnd=1
zu_bcnd=1
# path to output directory
outdir=./

# Enable float32 output
# Enable float32 density field
n_out_float32=1
out_float32_density=1

# Uncomment this to enable momentum_x
# out_float32_momentum_x=1

#################################################
# Parameters for linear wave problems
# initial density
rho=1.0
# velocity in the x direction
vx=0
# velocity in the y direction
vy=0
# velocity in the z direction
vz=0
# initial pressure
P=0.6
# amplitude of perturbing oscillations
A=1e-4
# value of gamma
gamma=1.666666666666667

26 changes: 25 additions & 1 deletion src/global/global.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -223,6 +223,30 @@ void parse_param(char *name,char *value, struct parameters *parms){
parms->n_rotated_projection = atoi(value);
else if (strcmp(name, "n_slice")==0)
parms->n_slice = atoi(value);
else if (strcmp(name, "n_out_float32")==0)
parms->n_out_float32 = atoi(value);
else if (strcmp(name, "out_float32_density")==0)
parms->out_float32_density = atoi(value);
else if (strcmp(name, "out_float32_momentum_x")==0)
parms->out_float32_momentum_x = atoi(value);
else if (strcmp(name, "out_float32_momentum_y")==0)
parms->out_float32_momentum_y = atoi(value);
else if (strcmp(name, "out_float32_momentum_z")==0)
parms->out_float32_momentum_z = atoi(value);
else if (strcmp(name, "out_float32_Energy")==0)
parms->out_float32_Energy = atoi(value);
#ifdef DE
else if (strcmp(name, "out_float32_GasEnergy")==0)
parms->out_float32_GasEnergy = atoi(value);
#endif // DE
#ifdef MHD
else if (strcmp(name, "out_float32_magnetic_x")==0)
parms->out_float32_magnetic_x = atoi(value);
else if (strcmp(name, "out_float32_magnetic_y")==0)
parms->out_float32_magnetic_y = atoi(value);
else if (strcmp(name, "out_float32_magnetic_z")==0)
parms->out_float32_magnetic_z = atoi(value);
#endif // MHD
else if (strcmp(name, "xmin")==0)
parms->xmin = atof(value);
else if (strcmp(name, "ymin")==0)
Expand Down Expand Up @@ -366,7 +390,7 @@ void parse_param(char *name,char *value, struct parameters *parms){
#ifdef CHEMISTRY_GPU
else if (strcmp(name, "UVB_rates_file")==0)
strncpy (parms->UVB_rates_file, value, MAXLEN);
#endif
#endif
#ifdef COOLING_GRACKLE
else if (strcmp(name, "UVB_rates_file")==0)
strncpy (parms->UVB_rates_file, value, MAXLEN);
Expand Down
14 changes: 14 additions & 0 deletions src/global/global.h
Original file line number Diff line number Diff line change
Expand Up @@ -204,6 +204,20 @@ struct parameters
int n_projection;
int n_rotated_projection;
int n_slice;
int n_out_float32=0;
int out_float32_density=0;
int out_float32_momentum_x=0;
int out_float32_momentum_y=0;
int out_float32_momentum_z=0;
int out_float32_Energy=0;
#ifdef DE
int out_float32_GasEnergy=0;
#endif
#ifdef MHD
int out_float32_magnetic_x=0;
int out_float32_magnetic_y=0;
int out_float32_magnetic_z=0;
#endif
Real xmin;
Real ymin;
Real zmin;
Expand Down
16 changes: 0 additions & 16 deletions src/global/global_cuda.h
Original file line number Diff line number Diff line change
Expand Up @@ -92,17 +92,6 @@ inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
}
}



/*! \fn Real minof3(Real a, Real b, Real c)
* \brief Returns the minimum of three floating point numbers. */
__device__ inline Real minof3(Real a, Real b, Real c)
{
return fmin(a, fmin(b,c));
}



/*! \fn int sgn_CUDA
* \brief Mathematical sign function. Returns sign of x. */
__device__ inline int sgn_CUDA(Real x)
Expand All @@ -111,11 +100,6 @@ __device__ inline int sgn_CUDA(Real x)
else return 1;
}


__global__ void test_function();



#endif //GLOBAL_CUDA_H

#endif //CUDA
18 changes: 0 additions & 18 deletions src/gravity/gravity_cuda.h

This file was deleted.

15 changes: 7 additions & 8 deletions src/gravity/gravity_cuda.cu → src/gravity/static_grav.h
Original file line number Diff line number Diff line change
Expand Up @@ -4,17 +4,16 @@
functions in hydro_cuda.cu. */
#ifdef CUDA

#pragma once

#include <stdio.h>
#include <math.h>
#include "../utils/gpu.hpp"
#include "../global/global.h"
#include "../global/global_cuda.h"
#include "../gravity/gravity_cuda.h"
#include <math.h> // provides sqrt log cos sin atan etc.
#include "../global/global.h" // provides GN etc.

// Work around lack of pow(Real,int) in Hip Clang for Rocm 3.5
static inline __device__ Real pow2(const Real x) { return x*x; }

__device__ void calc_g_1D(int xid, int x_off, int n_ghost, Real dx, Real xbound, Real *gx)
inline __device__ void calc_g_1D(int xid, int x_off, int n_ghost, Real dx, Real xbound, Real *gx)
{
Real x_pos, r_disk, r_halo;
x_pos = (x_off + xid - n_ghost + 0.5)*dx + xbound;
Expand Down Expand Up @@ -52,7 +51,7 @@ __device__ void calc_g_1D(int xid, int x_off, int n_ghost, Real dx, Real xbound,
}


__device__ void calc_g_2D(int xid, int yid, int x_off, int y_off, int n_ghost, Real dx, Real dy, Real xbound, Real ybound, Real *gx, Real *gy)
inline __device__ void calc_g_2D(int xid, int yid, int x_off, int y_off, int n_ghost, Real dx, Real dy, Real xbound, Real ybound, Real *gx, Real *gy)
{
Real x_pos, y_pos, r, phi;
// use the subgrid offset and global boundaries to calculate absolute positions on the grid
Expand Down Expand Up @@ -108,7 +107,7 @@ __device__ void calc_g_2D(int xid, int yid, int x_off, int y_off, int n_ghost, R
}


__device__ void calc_g_3D(int xid, int yid, int zid, 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 *gx, Real *gy, Real *gz)
inline __device__ void calc_g_3D(int xid, int yid, int zid, 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 *gx, Real *gy, Real *gz)
{
Real x_pos, y_pos, z_pos, r_disk, r_halo;
// use the subgrid offset and global boundaries to calculate absolute positions on the grid
Expand Down
2 changes: 1 addition & 1 deletion src/hydro/hydro_cuda.cu
Original file line number Diff line number Diff line change
Expand Up @@ -10,7 +10,7 @@
#include "../global/global.h"
#include "../global/global_cuda.h"
#include "../hydro/hydro_cuda.h"
#include "../gravity/gravity_cuda.h"
#include "../gravity/static_grav.h"
#include "../utils/hydro_utilities.h"
#include "../utils/cuda_utilities.h"
#include "../utils/reduction_utilities.h"
Expand Down
Loading

0 comments on commit d2a851d

Please sign in to comment.