From b38692c6dca7bfe8d2ab51a0872b631c44eaf918 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Fri, 8 Dec 2023 15:48:06 -0500 Subject: [PATCH 1/6] refactor to set dust grain size in input file --- src/dust/dust_cuda.cu | 21 +++++++++++---------- src/dust/dust_cuda.h | 6 +++--- src/global/global.cpp | 6 ++++++ src/global/global.h | 5 +++++ src/grid/grid3D.cpp | 8 +++++++- src/grid/grid3D.h | 6 ++++++ 6 files changed, 38 insertions(+), 14 deletions(-) diff --git a/src/dust/dust_cuda.cu b/src/dust/dust_cuda.cu index bbecf1935..fe3676e65 100644 --- a/src/dust/dust_cuda.cu +++ b/src/dust/dust_cuda.cu @@ -25,17 +25,17 @@ #include "../utils/gpu.hpp" #include "../utils/hydro_utilities.h" -void Dust_Update(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real dt, Real gamma) +void Dust_Update(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real dt, Real gamma, Real grain_radius) { int n_cells = nx * ny * nz; int ngrid = (n_cells + TPB - 1) / TPB; dim3 dim1dGrid(ngrid, 1, 1); dim3 dim1dBlock(TPB, 1, 1); - hipLaunchKernelGGL(Dust_Kernel, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, n_fields, dt, gamma); + hipLaunchKernelGGL(Dust_Kernel, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, n_fields, dt, gamma, grain_radius); CudaCheckError(); } -__global__ void Dust_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real dt, Real gamma) +__global__ void Dust_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real dt, Real gamma, Real grain_radius) { // get grid indices int n_cells = nx * ny * nz; @@ -100,7 +100,7 @@ __global__ void Dust_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_g temperature = temperature_init; Real tau_sp = - Calc_Sputtering_Timescale(number_density, temperature) / TIME_UNIT; // sputtering timescale, kyr (sim units) + Calc_Sputtering_Timescale(number_density, temperature, grain_radius) / TIME_UNIT; // sputtering timescale, kyr (sim units) dd_dt = Calc_dd_dt(density_dust, tau_sp); // rate of change in dust density at current timestep dd = dd_dt * dt; // change in dust density at current timestep @@ -126,17 +126,18 @@ __global__ void Dust_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_g } // McKinnon et al. (2017) sputtering timescale -__device__ __host__ Real Calc_Sputtering_Timescale(Real number_density, Real temperature) +__device__ __host__ Real Calc_Sputtering_Timescale(Real number_density, Real temperature, Real grain_radius) { - Real grain_radius = 1; // dust grain size in units of 0.1 micrometers - Real temperature_0 = 2e6; // temp above which the sputtering rate is ~constant in K - Real omega = 2.5; // controls the low-temperature scaling of the sputtering rate - Real A = 5.3618e15; // 0.17 Gyr in s + Real a = grain_radius; // dust grain size in units of 0.1 micrometers + Real temperature_0 = 2e6; // temp above which the sputtering rate is ~constant in K + Real omega = 2.5; // controls the low-temperature scaling of the sputtering rate + Real A = 5.3618e15; // 0.17 Gyr in s number_density /= (6e-4); // gas number density in units of 10^-27 g/cm^3 // sputtering timescale, s - Real tau_sp = A * (grain_radius / number_density) * (pow(temperature_0 / temperature, omega) + 1); + printf("%e\n", grain_radius); + Real tau_sp = A * (a / number_density) * (pow(temperature_0 / temperature, omega) + 1); return tau_sp; } diff --git a/src/dust/dust_cuda.h b/src/dust/dust_cuda.h index fb72007ac..ff27e4098 100644 --- a/src/dust/dust_cuda.h +++ b/src/dust/dust_cuda.h @@ -27,7 +27,7 @@ * \param[in] dt Simulation timestep * \param[in] gamma Specific heat ratio */ -void Dust_Update(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real dt, Real gamma); +void Dust_Update(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real dt, Real gamma, Real grain_radius); /*! * \brief Compute the change in dust density for a cell and update its value in dev_conserved. @@ -43,7 +43,7 @@ void Dust_Update(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n * \param[in] gamma Specific heat ratio */ __global__ void Dust_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real dt, - Real gamma); + Real gamma, Real grain_radius); /*! * \brief Compute the sputtering timescale based on a cell's density and temperature. @@ -53,7 +53,7 @@ __global__ void Dust_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_g * * \return Real Sputtering timescale in seconds (McKinnon et al. 2017) */ -__device__ __host__ Real Calc_Sputtering_Timescale(Real number_density, Real temperature); +__device__ __host__ Real Calc_Sputtering_Timescale(Real number_density, Real temperature, Real grain_radius); /*! * \brief Compute the rate of change in dust density based on the current dust density and sputtering timescale. diff --git a/src/global/global.cpp b/src/global/global.cpp index ecb7f2ccb..f57f2d475 100644 --- a/src/global/global.cpp +++ b/src/global/global.cpp @@ -422,6 +422,12 @@ void Parse_Param(char *name, char *value, struct Parameters *parms) strncpy(parms->skewersdir, value, MAXLEN); #endif #endif + #ifdef SCALAR + #ifdef DUST + } else if (strcmp(name, "grain_radius") == 0) { + parms->grain_radius = atoi(value); + #endif + #endif } else if (!Is_Param_Valid(name)) { chprintf("WARNING: %s/%s: Unknown parameter/value pair!\n", name, value); } diff --git a/src/global/global.h b/src/global/global.h index 8abe358fc..4e9148e79 100644 --- a/src/global/global.h +++ b/src/global/global.h @@ -307,6 +307,11 @@ struct Parameters { char skewersdir[MAXLEN]; #endif #endif +#ifdef SCALAR +#ifdef DUST +Real grain_radius; +#endif +#endif }; /*! \fn void parse_params(char *param_file, struct Parameters * parms); diff --git a/src/grid/grid3D.cpp b/src/grid/grid3D.cpp index 8a7d3fa6f..b7e44e2fe 100644 --- a/src/grid/grid3D.cpp +++ b/src/grid/grid3D.cpp @@ -274,6 +274,12 @@ void Grid3D::Initialize(struct Parameters *P) H.OUTPUT_SCALE_FACOR = not(P->scale_outputs_file[0] == '\0'); #endif +#ifdef SCALAR +#ifdef DUST + H.grain_radius = P->grain_radius; +#endif +#endif + H.Output_Initial = true; } @@ -494,7 +500,7 @@ Real Grid3D::Update_Grid(void) #ifdef DUST // ==Apply dust from dust/dust_cuda.h== - Dust_Update(C.device, H.nx, H.ny, H.nz, H.n_ghost, H.n_fields, H.dt, gama); + Dust_Update(C.device, H.nx, H.ny, H.nz, H.n_ghost, H.n_fields, H.dt, gama, H.grain_radius); #endif // DUST // Update the H and He ionization fractions and apply cooling and photoheating diff --git a/src/grid/grid3D.h b/src/grid/grid3D.h index 3f4d4772c..34bd57f7f 100644 --- a/src/grid/grid3D.h +++ b/src/grid/grid3D.h @@ -267,6 +267,12 @@ struct Header { * \brief Flag set to true when all the data will be written to file * (Restart File ) */ bool Output_Complete_Data; + + #ifdef SCALAR + #ifdef DUST + Real grain_radius; + #endif + #endif }; /*! \class Grid3D From 69d586c04c4266e1f23c26a693d2a1d784a3bdee Mon Sep 17 00:00:00 2001 From: helenarichie Date: Fri, 8 Dec 2023 15:58:20 -0500 Subject: [PATCH 2/6] wrap dust macro in scalar macro --- src/grid/initial_conditions.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/src/grid/initial_conditions.cpp b/src/grid/initial_conditions.cpp index 768bf0960..b363f51e9 100644 --- a/src/grid/initial_conditions.cpp +++ b/src/grid/initial_conditions.cpp @@ -1415,9 +1415,11 @@ void Grid3D::Clouds() C.GasEnergy[id] = p_cl / (gama - 1.0); #endif // DE -#ifdef DUST +#ifdef SCALAR + #ifdef DUST C.host[id + H.n_cells * grid_enum::dust_density] = rho_cl * 1e-2; -#endif // DUST + #endif // DUST +#endif } } } From 61970fca9936a3bf32cf56bd88a8fb734eed41a3 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Fri, 8 Dec 2023 16:02:06 -0500 Subject: [PATCH 3/6] run clang format --- src/dust/dust_cuda.cu | 13 ++++++++----- src/dust/dust_cuda.h | 7 ++++--- src/global/global.cpp | 6 +++--- src/global/global.h | 6 +++--- src/grid/grid3D.cpp | 4 ++-- src/grid/grid3D.h | 6 +++--- 6 files changed, 23 insertions(+), 19 deletions(-) diff --git a/src/dust/dust_cuda.cu b/src/dust/dust_cuda.cu index fe3676e65..40ae32236 100644 --- a/src/dust/dust_cuda.cu +++ b/src/dust/dust_cuda.cu @@ -25,17 +25,20 @@ #include "../utils/gpu.hpp" #include "../utils/hydro_utilities.h" -void Dust_Update(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real dt, Real gamma, Real grain_radius) +void Dust_Update(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real dt, Real gamma, + Real grain_radius) { int n_cells = nx * ny * nz; int ngrid = (n_cells + TPB - 1) / TPB; dim3 dim1dGrid(ngrid, 1, 1); dim3 dim1dBlock(TPB, 1, 1); - hipLaunchKernelGGL(Dust_Kernel, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, n_fields, dt, gamma, grain_radius); + hipLaunchKernelGGL(Dust_Kernel, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, n_fields, dt, gamma, + grain_radius); CudaCheckError(); } -__global__ void Dust_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real dt, Real gamma, Real grain_radius) +__global__ void Dust_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real dt, Real gamma, + Real grain_radius) { // get grid indices int n_cells = nx * ny * nz; @@ -99,8 +102,8 @@ __global__ void Dust_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_g // if dual energy is turned on use temp from total internal energy temperature = temperature_init; - Real tau_sp = - Calc_Sputtering_Timescale(number_density, temperature, grain_radius) / TIME_UNIT; // sputtering timescale, kyr (sim units) + Real tau_sp = Calc_Sputtering_Timescale(number_density, temperature, grain_radius) / + TIME_UNIT; // sputtering timescale, kyr (sim units) dd_dt = Calc_dd_dt(density_dust, tau_sp); // rate of change in dust density at current timestep dd = dd_dt * dt; // change in dust density at current timestep diff --git a/src/dust/dust_cuda.h b/src/dust/dust_cuda.h index ff27e4098..212901e8a 100644 --- a/src/dust/dust_cuda.h +++ b/src/dust/dust_cuda.h @@ -27,7 +27,8 @@ * \param[in] dt Simulation timestep * \param[in] gamma Specific heat ratio */ -void Dust_Update(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real dt, Real gamma, Real grain_radius); +void Dust_Update(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real dt, Real gamma, + Real grain_radius); /*! * \brief Compute the change in dust density for a cell and update its value in dev_conserved. @@ -42,8 +43,8 @@ void Dust_Update(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n * \param[in] dt Simulation timestep * \param[in] gamma Specific heat ratio */ -__global__ void Dust_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real dt, - Real gamma, Real grain_radius); +__global__ void Dust_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real dt, Real gamma, + Real grain_radius); /*! * \brief Compute the sputtering timescale based on a cell's density and temperature. diff --git a/src/global/global.cpp b/src/global/global.cpp index f57f2d475..9f828fd56 100644 --- a/src/global/global.cpp +++ b/src/global/global.cpp @@ -422,12 +422,12 @@ void Parse_Param(char *name, char *value, struct Parameters *parms) strncpy(parms->skewersdir, value, MAXLEN); #endif #endif - #ifdef SCALAR - #ifdef DUST +#ifdef SCALAR + #ifdef DUST } else if (strcmp(name, "grain_radius") == 0) { parms->grain_radius = atoi(value); - #endif #endif +#endif } else if (!Is_Param_Valid(name)) { chprintf("WARNING: %s/%s: Unknown parameter/value pair!\n", name, value); } diff --git a/src/global/global.h b/src/global/global.h index 4e9148e79..83aabd5d6 100644 --- a/src/global/global.h +++ b/src/global/global.h @@ -308,9 +308,9 @@ struct Parameters { #endif #endif #ifdef SCALAR -#ifdef DUST -Real grain_radius; -#endif + #ifdef DUST + Real grain_radius; + #endif #endif }; diff --git a/src/grid/grid3D.cpp b/src/grid/grid3D.cpp index b7e44e2fe..38384af84 100644 --- a/src/grid/grid3D.cpp +++ b/src/grid/grid3D.cpp @@ -275,9 +275,9 @@ void Grid3D::Initialize(struct Parameters *P) #endif #ifdef SCALAR -#ifdef DUST + #ifdef DUST H.grain_radius = P->grain_radius; -#endif + #endif #endif H.Output_Initial = true; diff --git a/src/grid/grid3D.h b/src/grid/grid3D.h index 34bd57f7f..4398dfe44 100644 --- a/src/grid/grid3D.h +++ b/src/grid/grid3D.h @@ -268,11 +268,11 @@ struct Header { * (Restart File ) */ bool Output_Complete_Data; - #ifdef SCALAR +#ifdef SCALAR #ifdef DUST - Real grain_radius; - #endif + Real grain_radius; #endif +#endif }; /*! \class Grid3D From 38b16493a06d7fd98ddc706569804bcd936a83ed Mon Sep 17 00:00:00 2001 From: helenarichie Date: Fri, 8 Dec 2023 16:24:56 -0500 Subject: [PATCH 4/6] undo accidental commit --- src/dust/dust_cuda.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/dust/dust_cuda.cu b/src/dust/dust_cuda.cu index 40ae32236..46273ef03 100644 --- a/src/dust/dust_cuda.cu +++ b/src/dust/dust_cuda.cu @@ -34,7 +34,7 @@ void Dust_Update(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n dim3 dim1dBlock(TPB, 1, 1); hipLaunchKernelGGL(Dust_Kernel, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, n_fields, dt, gamma, grain_radius); - CudaCheckError(); + GPU_Error_Check(); } __global__ void Dust_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real dt, Real gamma, From 7a02d014964dfe042afe13bbacbbb0e6360b412c Mon Sep 17 00:00:00 2001 From: helenarichie Date: Fri, 8 Dec 2023 16:30:24 -0500 Subject: [PATCH 5/6] update test --- src/dust/dust_cuda_tests.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/dust/dust_cuda_tests.cpp b/src/dust/dust_cuda_tests.cpp index a1357037a..4beab7df5 100644 --- a/src/dust/dust_cuda_tests.cpp +++ b/src/dust/dust_cuda_tests.cpp @@ -28,9 +28,10 @@ TEST(tDUSTTestSputteringTimescale, Real YR_IN_S = 3.154e7; Real const k_test_number_density = 1; Real const k_test_temperature = pow(10, 5.0); + Real const k_test_grain_size = 1; Real const k_fiducial_num = 182565146.96398282; - Real test_num = Calc_Sputtering_Timescale(k_test_number_density, k_test_temperature) / YR_IN_S; // yr + Real test_num = Calc_Sputtering_Timescale(k_test_number_density, k_test_temperature, k_test_grain_size) / YR_IN_S; // yr double abs_diff; int64_t ulps_diff; From 479051a0b452810bfb5fd6983c80c696110d3448 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Fri, 8 Dec 2023 16:32:00 -0500 Subject: [PATCH 6/6] run clang format --- src/dust/dust_cuda_tests.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/src/dust/dust_cuda_tests.cpp b/src/dust/dust_cuda_tests.cpp index 4beab7df5..fb0677edf 100644 --- a/src/dust/dust_cuda_tests.cpp +++ b/src/dust/dust_cuda_tests.cpp @@ -28,10 +28,11 @@ TEST(tDUSTTestSputteringTimescale, Real YR_IN_S = 3.154e7; Real const k_test_number_density = 1; Real const k_test_temperature = pow(10, 5.0); - Real const k_test_grain_size = 1; + Real const k_test_grain_radius = 1; Real const k_fiducial_num = 182565146.96398282; - Real test_num = Calc_Sputtering_Timescale(k_test_number_density, k_test_temperature, k_test_grain_size) / YR_IN_S; // yr + Real test_num = + Calc_Sputtering_Timescale(k_test_number_density, k_test_temperature, k_test_grain_radius) / YR_IN_S; // yr double abs_diff; int64_t ulps_diff;