From 3ecb87dce89e7714ea0c9ce4fc05373bec0b2d59 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Fri, 28 Apr 2023 16:39:37 -0400 Subject: [PATCH 01/40] add function to apply a floor to any conserved variable and apply a dust density floor --- src/hydro/hydro_cuda.cu | 26 ++++++++++++++++++++++++++ src/hydro/hydro_cuda.h | 5 +++++ src/integrators/VL_3D_cuda.cu | 6 ++++++ src/integrators/simple_3D_cuda.cu | 6 ++++++ 4 files changed, 43 insertions(+) diff --git a/src/hydro/hydro_cuda.cu b/src/hydro/hydro_cuda.cu index dad6f3b66..1cafca97f 100644 --- a/src/hydro/hydro_cuda.cu +++ b/src/hydro/hydro_cuda.cu @@ -1170,4 +1170,30 @@ __device__ void Average_Cell_All_Fields(int i, int j, int k, int nx, int ny, int #endif // DE } + #ifdef CONSERVED_FLOOR +__global__ void Apply_Conserved_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, + Real conserved_floor) +{ + int id, xid, yid, zid, n_cells; + Real field_0; + n_cells = nx * ny * nz; + + // get a global thread ID + id = threadIdx.x + blockIdx.x * blockDim.x; + zid = id / (nx * ny); + yid = (id - zid * nx * ny) / nx; + xid = id - zid * nx * ny - yid * nx; + + // threads corresponding to real cells do the calculation + 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) { + field_0 = dev_conserved[id + n_cells * field_num]; + + if (field_0 < conserved_floor) { + dev_conserved[id + n_cells * field_num] = conserved_floor; + } + } +} + #endif // CONSERVED_FLOOR + #endif // CUDA diff --git a/src/hydro/hydro_cuda.h b/src/hydro/hydro_cuda.h index a5c4ab713..bbd8ba5db 100644 --- a/src/hydro/hydro_cuda.h +++ b/src/hydro/hydro_cuda.h @@ -89,6 +89,11 @@ __global__ void Apply_Temperature_Floor(Real *dev_conserved, int nx, int ny, int Real U_floor); #endif + #ifdef CONSERVED_FLOOR +__global__ void Apply_Conserved_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, + Real conserved_floor); + #endif + __global__ void Partial_Update_Advected_Internal_Energy_1D(Real *dev_conserved, Real *Q_Lx, Real *Q_Rx, int nx, int n_ghost, Real dx, Real dt, Real gamma, int n_fields); diff --git a/src/integrators/VL_3D_cuda.cu b/src/integrators/VL_3D_cuda.cu index 2227172bf..cf5c5d4b8 100644 --- a/src/integrators/VL_3D_cuda.cu +++ b/src/integrators/VL_3D_cuda.cu @@ -328,6 +328,12 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int CudaCheckError(); #endif // TEMPERATURE_FLOOR + #ifdef DUST_FLOOR + hipLaunchKernelGGL(Apply_Conserved_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, + grid_enum::dust_density, 1e-5); + CudaCheckError(); + #endif // DUST_FLOOR + return; } diff --git a/src/integrators/simple_3D_cuda.cu b/src/integrators/simple_3D_cuda.cu index 01c9c6ac1..694eddac3 100644 --- a/src/integrators/simple_3D_cuda.cu +++ b/src/integrators/simple_3D_cuda.cu @@ -188,6 +188,12 @@ void Simple_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, CudaCheckError(); #endif // TEMPERATURE_FLOOR + #ifdef DUST_FLOOR + hipLaunchKernelGGL(Apply_Conserved_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, + grid_enum::dust_density, 1e-5); + CudaCheckError(); + #endif // DUST_FLOOR + return; } From 3dd1d300856e013ebe12d6d84070b995c872922f Mon Sep 17 00:00:00 2001 From: helenarichie Date: Fri, 28 Apr 2023 16:46:39 -0400 Subject: [PATCH 02/40] remove unnecessary macro --- src/hydro/hydro_cuda.cu | 2 -- src/hydro/hydro_cuda.h | 2 -- src/integrators/VL_3D_cuda.cu | 4 ++-- src/integrators/simple_3D_cuda.cu | 4 ++-- 4 files changed, 4 insertions(+), 8 deletions(-) diff --git a/src/hydro/hydro_cuda.cu b/src/hydro/hydro_cuda.cu index 1cafca97f..f53836fb7 100644 --- a/src/hydro/hydro_cuda.cu +++ b/src/hydro/hydro_cuda.cu @@ -1170,7 +1170,6 @@ __device__ void Average_Cell_All_Fields(int i, int j, int k, int nx, int ny, int #endif // DE } - #ifdef CONSERVED_FLOOR __global__ void Apply_Conserved_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, Real conserved_floor) { @@ -1194,6 +1193,5 @@ __global__ void Apply_Conserved_Floor(Real *dev_conserved, int nx, int ny, int n } } } - #endif // CONSERVED_FLOOR #endif // CUDA diff --git a/src/hydro/hydro_cuda.h b/src/hydro/hydro_cuda.h index bbd8ba5db..287224c57 100644 --- a/src/hydro/hydro_cuda.h +++ b/src/hydro/hydro_cuda.h @@ -89,10 +89,8 @@ __global__ void Apply_Temperature_Floor(Real *dev_conserved, int nx, int ny, int Real U_floor); #endif - #ifdef CONSERVED_FLOOR __global__ void Apply_Conserved_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, Real conserved_floor); - #endif __global__ void Partial_Update_Advected_Internal_Energy_1D(Real *dev_conserved, Real *Q_Lx, Real *Q_Rx, int nx, int n_ghost, Real dx, Real dt, Real gamma, int n_fields); diff --git a/src/integrators/VL_3D_cuda.cu b/src/integrators/VL_3D_cuda.cu index cf5c5d4b8..3fed3a9bf 100644 --- a/src/integrators/VL_3D_cuda.cu +++ b/src/integrators/VL_3D_cuda.cu @@ -328,11 +328,11 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int CudaCheckError(); #endif // TEMPERATURE_FLOOR - #ifdef DUST_FLOOR + #ifdef DUST hipLaunchKernelGGL(Apply_Conserved_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, grid_enum::dust_density, 1e-5); CudaCheckError(); - #endif // DUST_FLOOR + #endif // DUST return; } diff --git a/src/integrators/simple_3D_cuda.cu b/src/integrators/simple_3D_cuda.cu index 694eddac3..4a24bb362 100644 --- a/src/integrators/simple_3D_cuda.cu +++ b/src/integrators/simple_3D_cuda.cu @@ -188,11 +188,11 @@ void Simple_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, CudaCheckError(); #endif // TEMPERATURE_FLOOR - #ifdef DUST_FLOOR + #ifdef DUST hipLaunchKernelGGL(Apply_Conserved_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, grid_enum::dust_density, 1e-5); CudaCheckError(); - #endif // DUST_FLOOR + #endif // DUST return; } From 2a6eb5fab84679659701f57061dd1a52c837fb56 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Tue, 15 Aug 2023 10:24:55 -0400 Subject: [PATCH 03/40] save progress --- builds/make.type.dust | 1 + cholla-tests-data | 2 +- src/hydro/hydro_cuda.cu | 48 +++++++++++++++++++++++++------ src/hydro/hydro_cuda.h | 8 +++--- src/integrators/VL_3D_cuda.cu | 39 +++++++------------------ src/integrators/simple_3D_cuda.cu | 6 ++-- 6 files changed, 59 insertions(+), 45 deletions(-) diff --git a/builds/make.type.dust b/builds/make.type.dust index 0be259763..6570f7d8a 100644 --- a/builds/make.type.dust +++ b/builds/make.type.dust @@ -16,6 +16,7 @@ DFLAGS += -DHLLC # DFLAGS += -DDE DFLAGS += -DAVERAGE_SLOW_CELLS DFLAGS += -DTEMPERATURE_FLOOR +DFLAGS += -DSCALAR_FLOOR ifeq ($(findstring cosmology,$(TYPE)),cosmology) DFLAGS += -DSIMPLE diff --git a/cholla-tests-data b/cholla-tests-data index dcd73ff52..321416680 160000 --- a/cholla-tests-data +++ b/cholla-tests-data @@ -1 +1 @@ -Subproject commit dcd73ff52b9027627b247c6d888bcdb56840c85e +Subproject commit 321416680f95d97b5d4ccc6f0b83a8b9ecafdaf0 diff --git a/src/hydro/hydro_cuda.cu b/src/hydro/hydro_cuda.cu index 1f4b91fd3..0a3ed342f 100644 --- a/src/hydro/hydro_cuda.cu +++ b/src/hydro/hydro_cuda.cu @@ -276,7 +276,7 @@ __global__ void Update_Conserved_Variables_3D(Real *dev_conserved, Real *Q_Lx, R // issues #endif - #ifdef DENSITY_FLOOR + #ifdef DENSITY_FLOOR if (dev_conserved[id] < density_floor) { if (dev_conserved[id] > 0) { dens_0 = dev_conserved[id]; @@ -1086,7 +1086,6 @@ __global__ void Sync_Energies_3D(Real *dev_conserved, int nx, int ny, int nz, in #endif // DE - #ifdef TEMPERATURE_FLOOR __global__ void Apply_Temperature_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real U_floor) { @@ -1124,8 +1123,39 @@ __global__ void Apply_Temperature_Floor(Real *dev_conserved, int nx, int ny, int #endif } } - #endif // TEMPERATURE_FLOOR +__global__ void Apply_Density_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, Real density_floor) +{ + int id, xid, yid, zid, n_cells; + Real density_init; // variable to store the value of the scalar before a floor is applied + n_cells = nx * ny * nz; + + // get a global thread ID + id = threadIdx.x + blockIdx.x * blockDim.x; + zid = id / (nx * ny); + yid = (id - zid * nx * ny) / nx; + xid = id - zid * nx * ny - yid * nx; + + // threads corresponding to real cells do the calculation + 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) { + + density_init = dev_conserved[id + n_cells * grid_enum::density]; + + if (density_init < density_floor) { + printf("###Thread density change %f -> %f \n", density_init, density_floor); + dev_conserved[id] = density_floor; + // Scale the conserved values to the new density + dev_conserved[id + n_cells * grid_enum::momentum_x] *= (density_floor / density_init); + dev_conserved[id + n_cells * grid_enum::momentum_y] *= (density_floor / density_init); + dev_conserved[id + n_cells * grid_enum::momentum_z] *= (density_floor / density_init); + dev_conserved[id + n_cells * grid_enum::Energy] *= (density_floor / density_init); + #ifdef DE + dev_conserved[id + n_cells * grid_enum::GasEnergy] *= (density_floor / density_init); + #endif // DE + } + } +} __device__ Real Average_Cell_Single_Field(int field_indx, int i, int j, int k, int nx, int ny, int nz, int ncells, Real *conserved) { @@ -1170,11 +1200,11 @@ __device__ void Average_Cell_All_Fields(int i, int j, int k, int nx, int ny, int #endif // DE } -__global__ void Apply_Conserved_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, - Real conserved_floor) +__global__ void Apply_Scalar_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, + Real scalar_floor) { int id, xid, yid, zid, n_cells; - Real field_0; + Real scalar; // variable to store the value of the scalar before a floor is applied n_cells = nx * ny * nz; // get a global thread ID @@ -1186,10 +1216,10 @@ __global__ void Apply_Conserved_Floor(Real *dev_conserved, int nx, int ny, int n // threads corresponding to real cells do the calculation 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) { - field_0 = dev_conserved[id + n_cells * field_num]; + scalar = dev_conserved[id + n_cells * field_num]; - if (field_0 < conserved_floor) { - dev_conserved[id + n_cells * field_num] = conserved_floor; + if (scalar < scalar_floor) { + dev_conserved[id + n_cells * field_num] = scalar_floor; } } } diff --git a/src/hydro/hydro_cuda.h b/src/hydro/hydro_cuda.h index 287224c57..1f7d9a473 100644 --- a/src/hydro/hydro_cuda.h +++ b/src/hydro/hydro_cuda.h @@ -84,13 +84,13 @@ __global__ void Average_Slow_Cells_3D(Real *dev_conserved, int nx, int ny, int n Real dy, Real dz, Real gamma, Real max_dti_slow); #endif - #ifdef TEMPERATURE_FLOOR __global__ void Apply_Temperature_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real U_floor); - #endif -__global__ void Apply_Conserved_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, - Real conserved_floor); +__global__ void Apply_Density_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, Real density_floor); + +__global__ void Apply_Scalar_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, + Real scalar_floor); __global__ void Partial_Update_Advected_Internal_Energy_1D(Real *dev_conserved, Real *Q_Lx, Real *Q_Rx, int nx, int n_ghost, Real dx, Real dt, Real gamma, int n_fields); diff --git a/src/integrators/VL_3D_cuda.cu b/src/integrators/VL_3D_cuda.cu index 6ccc814c2..dd088d8ef 100644 --- a/src/integrators/VL_3D_cuda.cu +++ b/src/integrators/VL_3D_cuda.cu @@ -32,8 +32,7 @@ __global__ void Update_Conserved_Variables_3D_half(Real *dev_conserved, Real *dev_conserved_half, Real *dev_F_x, Real *dev_F_y, Real *dev_F_z, int nx, int ny, int nz, int n_ghost, - Real dx, Real dy, Real dz, Real dt, Real gamma, int n_fields, - Real density_floor); + Real dx, Real dy, Real dz, Real dt, Real gamma, int n_fields); 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, @@ -194,8 +193,13 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int // Step 3: Update the conserved variables half a timestep hipLaunchKernelGGL(Update_Conserved_Variables_3D_half, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, dev_conserved_half, - F_x, F_y, F_z, nx, ny, nz, n_ghost, dx, dy, dz, 0.5 * dt, gama, n_fields, density_floor); + F_x, F_y, F_z, nx, ny, nz, n_ghost, dx, dy, dz, 0.5 * dt, gama, n_fields); CudaCheckError(); + + #ifdef DENSITY_FLOOR + hipLaunchKernelGGL(Apply_Density_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved_half, nx, ny, nz, n_ghost, density_floor); + #endif // DENSITY_FLOOR + #ifdef MHD // Update the magnetic fields hipLaunchKernelGGL(mhd::Update_Magnetic_Field_3D, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, dev_conserved_half, @@ -325,11 +329,11 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int CudaCheckError(); #endif // TEMPERATURE_FLOOR - #ifdef DUST - hipLaunchKernelGGL(Apply_Conserved_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, + #ifdef SCALAR_FLOOR + hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, grid_enum::dust_density, 1e-5); CudaCheckError(); - #endif // DUST + #endif // SCALAR_FLOOR return; } @@ -353,8 +357,7 @@ void Free_Memory_VL_3D() __global__ void Update_Conserved_Variables_3D_half(Real *dev_conserved, Real *dev_conserved_half, Real *dev_F_x, Real *dev_F_y, Real *dev_F_z, int nx, int ny, int nz, int n_ghost, - Real dx, Real dy, Real dz, Real dt, Real gamma, int n_fields, - Real density_floor) + Real dx, Real dy, Real dz, Real dt, Real gamma, int n_fields) { Real dtodx = dt / dx; Real dtody = dt / dy; @@ -378,10 +381,6 @@ __global__ void Update_Conserved_Variables_3D_half(Real *dev_conserved, Real *de int ipo, jpo, kpo; #endif // DE - #ifdef DENSITY_FLOOR - Real dens_0; - #endif // DENSITY_FLOOR - // threads corresponding to all cells except outer ring of ghost cells do the // calculation if (xid > 0 && xid < nx - 1 && yid > 0 && yid < ny - 1 && zid > 0 && zid < nz - 1) { @@ -455,22 +454,6 @@ __global__ void Update_Conserved_Variables_3D_half(Real *dev_conserved, Real *de dtodz * (dev_F_z[(n_fields - 1) * n_cells + kmo] - dev_F_z[(n_fields - 1) * n_cells + id]) + 0.5 * P * (dtodx * (vx_imo - vx_ipo) + dtody * (vy_jmo - vy_jpo) + dtodz * (vz_kmo - vz_kpo)); #endif // DE - - #ifdef DENSITY_FLOOR - if (dev_conserved_half[id] < density_floor) { - dens_0 = dev_conserved_half[id]; - printf("###Thread density change %f -> %f \n", dens_0, density_floor); - dev_conserved_half[id] = density_floor; - // Scale the conserved values to the new density - dev_conserved_half[1 * n_cells + id] *= (density_floor / dens_0); - dev_conserved_half[2 * n_cells + id] *= (density_floor / dens_0); - dev_conserved_half[3 * n_cells + id] *= (density_floor / dens_0); - dev_conserved_half[4 * n_cells + id] *= (density_floor / dens_0); - #ifdef DE - dev_conserved_half[(n_fields - 1) * n_cells + id] *= (density_floor / dens_0); - #endif // DE - } - #endif // DENSITY_FLOOR } } diff --git a/src/integrators/simple_3D_cuda.cu b/src/integrators/simple_3D_cuda.cu index 3c8dc13e9..dfc015f25 100644 --- a/src/integrators/simple_3D_cuda.cu +++ b/src/integrators/simple_3D_cuda.cu @@ -185,11 +185,11 @@ void Simple_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, CudaCheckError(); #endif // TEMPERATURE_FLOOR - #ifdef DUST - hipLaunchKernelGGL(Apply_Conserved_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, + #ifdef SCALAR_FLOOR + hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, grid_enum::dust_density, 1e-5); CudaCheckError(); - #endif // DUST + #endif // SCALAR_FLOOR return; } From 77d6e1ab38573857951fb15ad1c608cccf591ea7 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Mon, 25 Sep 2023 15:39:56 -0400 Subject: [PATCH 04/40] add test for scalar floor kernel --- src/dust/dust_cuda_tests.cpp | 4 ++-- src/hydro/hydro_cuda_tests.cu | 45 +++++++++++++++++++++++++++++++++++ 2 files changed, 47 insertions(+), 2 deletions(-) diff --git a/src/dust/dust_cuda_tests.cpp b/src/dust/dust_cuda_tests.cpp index a1357037a..eeaac2368 100644 --- a/src/dust/dust_cuda_tests.cpp +++ b/src/dust/dust_cuda_tests.cpp @@ -22,7 +22,7 @@ #ifdef DUST TEST(tDUSTTestSputteringTimescale, - CorrectInputExpectCorrectOutput) // test suite name, test name + CorrectInputExpectCorrectOutput) { // Parameters Real YR_IN_S = 3.154e7; @@ -46,7 +46,7 @@ TEST(tDUSTTestSputteringTimescale, } TEST(tDUSTTestSputteringGrowthRate, - CorrectInputExpectCorrectOutput) // test suite name, test name + CorrectInputExpectCorrectOutput) { // Parameters Real YR_IN_S = 3.154e7; diff --git a/src/hydro/hydro_cuda_tests.cu b/src/hydro/hydro_cuda_tests.cu index fe4b351f7..8fd728528 100644 --- a/src/hydro/hydro_cuda_tests.cu +++ b/src/hydro/hydro_cuda_tests.cu @@ -146,4 +146,49 @@ TEST(tMHDMhdInverseCrossingTime, CorrectInputExpectCorrectOutput) // End of tests for the mhdInverseCrossingTime function // ============================================================================= +TEST(tHYDROScalarFloor, CorrectInputExpectCorrectOutput) +{ + // Call the function we are testing + int num_blocks = 1; + dim3 dim1dGrid(num_blocks, 1, 1); + dim3 dim1dBlock(TPB, 1, 1); + int const nx = 1; + int const ny = 1; + int const nz = 1; + int const n_fields = 6; // 5 conserved + 1 scalar + int const n_ghost = 0; + int const field_num = 5; // scalar field index + + // initialize host and device conserved arrays + std::vector host_conserved(n_fields); + cuda_utilities::DeviceVector dev_conserved(n_fields); + + // Set values of conserved variables for input (host) + host_conserved.at(0) = 0.0; // density + host_conserved.at(1) = 0.0; // x momentum + host_conserved.at(2) = 0.0; // y momentum + host_conserved.at(3) = 0.0; // z momentum + host_conserved.at(4) = 0.0; // energy + + Real scalar_floor = 1.0; // minimum allowed value for scalar field + + // Case where scalar is below the floor + host_conserved.at(field_num) = 0.0; // scalar + dev_conserved.cpyHostToDevice(host_conserved); + hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved.data(), nx, ny, nz, n_ghost, field_num, scalar_floor); + testing_utilities::Check_Results(scalar_floor, dev_conserved.at(field_num), "below floor"); + + // Case where scalar is above the floor + host_conserved.at(field_num) = 2.0; // scalar + dev_conserved.cpyHostToDevice(host_conserved); + hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved.data(), nx, ny, nz, n_ghost, field_num, scalar_floor); + testing_utilities::Check_Results(host_conserved.at(field_num), dev_conserved.at(field_num), "above floor"); + + // Case where scalar is at the floor + host_conserved.at(field_num) = 1.0; // scalar + dev_conserved.cpyHostToDevice(host_conserved); + hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved.data(), nx, ny, nz, n_ghost, field_num, scalar_floor); + testing_utilities::Check_Results(host_conserved.at(field_num), dev_conserved.at(field_num), "at floor"); +} + #endif // CUDA From 639be11a8369d604ea2351ac8f668b4d41e8ccc8 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Mon, 25 Sep 2023 15:44:10 -0400 Subject: [PATCH 05/40] run clang format --- src/dust/dust_cuda_tests.cpp | 6 ++---- src/hydro/hydro_cuda.cu | 17 ++++++++--------- src/hydro/hydro_cuda.h | 2 +- src/hydro/hydro_cuda_tests.cu | 25 ++++++++++++++----------- src/integrators/VL_3D_cuda.cu | 3 ++- 5 files changed, 27 insertions(+), 26 deletions(-) diff --git a/src/dust/dust_cuda_tests.cpp b/src/dust/dust_cuda_tests.cpp index eeaac2368..aa26f7a01 100644 --- a/src/dust/dust_cuda_tests.cpp +++ b/src/dust/dust_cuda_tests.cpp @@ -21,8 +21,7 @@ #ifdef DUST -TEST(tDUSTTestSputteringTimescale, - CorrectInputExpectCorrectOutput) +TEST(tDUSTTestSputteringTimescale, CorrectInputExpectCorrectOutput) { // Parameters Real YR_IN_S = 3.154e7; @@ -45,8 +44,7 @@ TEST(tDUSTTestSputteringTimescale, << "The ULP difference is: " << ulps_diff << std::endl; } -TEST(tDUSTTestSputteringGrowthRate, - CorrectInputExpectCorrectOutput) +TEST(tDUSTTestSputteringGrowthRate, CorrectInputExpectCorrectOutput) { // Parameters Real YR_IN_S = 3.154e7; diff --git a/src/hydro/hydro_cuda.cu b/src/hydro/hydro_cuda.cu index 0a3ed342f..037d41517 100644 --- a/src/hydro/hydro_cuda.cu +++ b/src/hydro/hydro_cuda.cu @@ -276,7 +276,7 @@ __global__ void Update_Conserved_Variables_3D(Real *dev_conserved, Real *Q_Lx, R // issues #endif - #ifdef DENSITY_FLOOR + #ifdef DENSITY_FLOOR if (dev_conserved[id] < density_floor) { if (dev_conserved[id] > 0) { dens_0 = dev_conserved[id]; @@ -1115,19 +1115,19 @@ __global__ void Apply_Temperature_Floor(Real *dev_conserved, int nx, int ny, int dev_conserved[4 * n_cells + id] = Ekin + d * U_floor; } - #ifdef DE + #ifdef DE U = dev_conserved[(n_fields - 1) * n_cells + id] / d; if (U < U_floor) { dev_conserved[(n_fields - 1) * n_cells + id] = d * U_floor; } - #endif + #endif } } __global__ void Apply_Density_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, Real density_floor) { int id, xid, yid, zid, n_cells; - Real density_init; // variable to store the value of the scalar before a floor is applied + Real density_init; // variable to store the value of the scalar before a floor is applied n_cells = nx * ny * nz; // get a global thread ID @@ -1136,10 +1136,9 @@ __global__ void Apply_Density_Floor(Real *dev_conserved, int nx, int ny, int nz, yid = (id - zid * nx * ny) / nx; xid = id - zid * nx * ny - yid * nx; - // threads corresponding to real cells do the calculation + // threads corresponding to real cells do the calculation 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) { - density_init = dev_conserved[id + n_cells * grid_enum::density]; if (density_init < density_floor) { @@ -1150,9 +1149,9 @@ __global__ void Apply_Density_Floor(Real *dev_conserved, int nx, int ny, int nz, dev_conserved[id + n_cells * grid_enum::momentum_y] *= (density_floor / density_init); dev_conserved[id + n_cells * grid_enum::momentum_z] *= (density_floor / density_init); dev_conserved[id + n_cells * grid_enum::Energy] *= (density_floor / density_init); - #ifdef DE + #ifdef DE dev_conserved[id + n_cells * grid_enum::GasEnergy] *= (density_floor / density_init); - #endif // DE + #endif // DE } } } @@ -1204,7 +1203,7 @@ __global__ void Apply_Scalar_Floor(Real *dev_conserved, int nx, int ny, int nz, Real scalar_floor) { int id, xid, yid, zid, n_cells; - Real scalar; // variable to store the value of the scalar before a floor is applied + Real scalar; // variable to store the value of the scalar before a floor is applied n_cells = nx * ny * nz; // get a global thread ID diff --git a/src/hydro/hydro_cuda.h b/src/hydro/hydro_cuda.h index 1f7d9a473..2e93248f8 100644 --- a/src/hydro/hydro_cuda.h +++ b/src/hydro/hydro_cuda.h @@ -90,7 +90,7 @@ __global__ void Apply_Temperature_Floor(Real *dev_conserved, int nx, int ny, int __global__ void Apply_Density_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, Real density_floor); __global__ void Apply_Scalar_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, - Real scalar_floor); + Real scalar_floor); __global__ void Partial_Update_Advected_Internal_Energy_1D(Real *dev_conserved, Real *Q_Lx, Real *Q_Rx, int nx, int n_ghost, Real dx, Real dt, Real gamma, int n_fields); diff --git a/src/hydro/hydro_cuda_tests.cu b/src/hydro/hydro_cuda_tests.cu index 8fd728528..da0f185e5 100644 --- a/src/hydro/hydro_cuda_tests.cu +++ b/src/hydro/hydro_cuda_tests.cu @@ -152,11 +152,11 @@ TEST(tHYDROScalarFloor, CorrectInputExpectCorrectOutput) int num_blocks = 1; dim3 dim1dGrid(num_blocks, 1, 1); dim3 dim1dBlock(TPB, 1, 1); - int const nx = 1; - int const ny = 1; - int const nz = 1; - int const n_fields = 6; // 5 conserved + 1 scalar - int const n_ghost = 0; + int const nx = 1; + int const ny = 1; + int const nz = 1; + int const n_fields = 6; // 5 conserved + 1 scalar + int const n_ghost = 0; int const field_num = 5; // scalar field index // initialize host and device conserved arrays @@ -174,20 +174,23 @@ TEST(tHYDROScalarFloor, CorrectInputExpectCorrectOutput) // Case where scalar is below the floor host_conserved.at(field_num) = 0.0; // scalar - dev_conserved.cpyHostToDevice(host_conserved); - hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved.data(), nx, ny, nz, n_ghost, field_num, scalar_floor); + dev_conserved.cpyHostToDevice(host_conserved); + hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved.data(), nx, ny, nz, n_ghost, + field_num, scalar_floor); testing_utilities::Check_Results(scalar_floor, dev_conserved.at(field_num), "below floor"); // Case where scalar is above the floor host_conserved.at(field_num) = 2.0; // scalar - dev_conserved.cpyHostToDevice(host_conserved); - hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved.data(), nx, ny, nz, n_ghost, field_num, scalar_floor); + dev_conserved.cpyHostToDevice(host_conserved); + hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved.data(), nx, ny, nz, n_ghost, + field_num, scalar_floor); testing_utilities::Check_Results(host_conserved.at(field_num), dev_conserved.at(field_num), "above floor"); // Case where scalar is at the floor host_conserved.at(field_num) = 1.0; // scalar - dev_conserved.cpyHostToDevice(host_conserved); - hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved.data(), nx, ny, nz, n_ghost, field_num, scalar_floor); + dev_conserved.cpyHostToDevice(host_conserved); + hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved.data(), nx, ny, nz, n_ghost, + field_num, scalar_floor); testing_utilities::Check_Results(host_conserved.at(field_num), dev_conserved.at(field_num), "at floor"); } diff --git a/src/integrators/VL_3D_cuda.cu b/src/integrators/VL_3D_cuda.cu index dd088d8ef..9b1fe5b86 100644 --- a/src/integrators/VL_3D_cuda.cu +++ b/src/integrators/VL_3D_cuda.cu @@ -197,7 +197,8 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int CudaCheckError(); #ifdef DENSITY_FLOOR - hipLaunchKernelGGL(Apply_Density_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved_half, nx, ny, nz, n_ghost, density_floor); + hipLaunchKernelGGL(Apply_Density_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved_half, nx, ny, nz, n_ghost, + density_floor); #endif // DENSITY_FLOOR #ifdef MHD From 0fd7f89d31b9325bfc4b626fe73c616117e71a64 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Wed, 6 Dec 2023 13:09:54 -0500 Subject: [PATCH 06/40] update branch --- builds/make.type.dust | 5 ++++- cholla-tests-data | 2 +- src/global/global.h | 2 +- src/grid/cuda_boundaries.cu | 4 ++-- src/grid/grid3D.cpp | 2 +- src/grid/initial_conditions.cpp | 15 ++++++++------- src/integrators/VL_3D_cuda.cu | 4 +++- src/integrators/simple_3D_cuda.cu | 2 ++ 8 files changed, 22 insertions(+), 14 deletions(-) diff --git a/builds/make.type.dust b/builds/make.type.dust index a0ea7b267..b9813674c 100644 --- a/builds/make.type.dust +++ b/builds/make.type.dust @@ -9,7 +9,7 @@ MPI_GPU ?= DFLAGS += -DCUDA DFLAGS += -DMPI_CHOLLA DFLAGS += -DPRECISION=2 -DFLAGS += -DPPMC +DFLAGS += -DPPMP DFLAGS += -DHLLC # DFLAGS += -DDE @@ -30,8 +30,11 @@ DFLAGS += -DSCALAR # Define dust macro DFLAGS += -DDUST +DFLAGS += -DSCALAR_FLOOR + # Apply the cooling in the GPU from precomputed tables DFLAGS += -DCOOLING_GPU +DFLAGS += -DCLOUDY_COOLING #Measure the Timing of the different stages #DFLAGS += -DCPU_TIME diff --git a/cholla-tests-data b/cholla-tests-data index 321416680..dcd73ff52 160000 --- a/cholla-tests-data +++ b/cholla-tests-data @@ -1 +1 @@ -Subproject commit 321416680f95d97b5d4ccc6f0b83a8b9ecafdaf0 +Subproject commit dcd73ff52b9027627b247c6d888bcdb56840c85e diff --git a/src/global/global.h b/src/global/global.h index 8abe358fc..1e63a6af7 100644 --- a/src/global/global.h +++ b/src/global/global.h @@ -50,7 +50,7 @@ typedef double Real; #define LOG_FILE_NAME "run_output.log" // Conserved Floor Values -#define TEMP_FLOOR 1e-3 +#define TEMP_FLOOR 10 #define DENS_FLOOR 1e-5 // in code units // Parameter for Enzo dual Energy Condition diff --git a/src/grid/cuda_boundaries.cu b/src/grid/cuda_boundaries.cu index f5dbe361d..733fb1a85 100644 --- a/src/grid/cuda_boundaries.cu +++ b/src/grid/cuda_boundaries.cu @@ -312,13 +312,13 @@ __global__ void Wind_Boundary_kernel(Real *c_device, int nx, int ny, int nz, int Real vx, vy, vz, d_0, P_0; n_0 = 1e-2; // same value as n_bg in cloud initial condition function (cm^-3) - T_0 = 3e6; // same value as T_bg in cloud initial condition function (K) + T_0 = 3e7; // same value as T_bg in cloud initial condition function (K) // same values as rho_bg and p_bg in cloud initial condition function d_0 = n_0 * mu * MP / DENSITY_UNIT; P_0 = n_0 * KB * T_0 / PRESSURE_UNIT; - vx = 100 * TIME_UNIT / KPC; // km/s * (cholla unit conversion) + vx = 500 * TIME_UNIT / KPC; // km/s * (cholla unit conversion) vy = 0.0; vz = 0.0; diff --git a/src/grid/grid3D.cpp b/src/grid/grid3D.cpp index 8a7d3fa6f..0a8374323 100644 --- a/src/grid/grid3D.cpp +++ b/src/grid/grid3D.cpp @@ -157,7 +157,7 @@ void Grid3D::Initialize(struct Parameters *P) C_cfl = 0.3; #ifdef AVERAGE_SLOW_CELLS - H.min_dt_slow = 1e-100; // Initialize the minumum dt to a tiny number + H.min_dt_slow = 1e-5; // Initialize the minumum dt to a tiny number #endif // AVERAGE_SLOW_CELLS #ifndef MPI_CHOLLA diff --git a/src/grid/initial_conditions.cpp b/src/grid/initial_conditions.cpp index 768bf0960..a1983fb48 100644 --- a/src/grid/initial_conditions.cpp +++ b/src/grid/initial_conditions.cpp @@ -1324,7 +1324,7 @@ void Grid3D::Clouds() Real p_bg, p_cl; // background and cloud pressure Real mu = 0.6; // mean atomic weight int N_cl = 1; // number of clouds - Real R_cl = 2.5; // cloud radius in code units (kpc) + Real R_cl = .1; // cloud radius in code units (kpc) Real cl_pos[N_cl][3]; // array of cloud positions Real r; @@ -1339,22 +1339,22 @@ void Grid3D::Clouds() // single centered cloud setup for (int nn = 0; nn < N_cl; nn++) { - cl_pos[nn][0] = 0.5 * H.xdglobal; + cl_pos[nn][0] = 0.075 * H.xdglobal; cl_pos[nn][1] = 0.5 * H.ydglobal; cl_pos[nn][2] = 0.5 * H.zdglobal; printf("Cloud positions: %f %f %f\n", cl_pos[nn][0], cl_pos[nn][1], cl_pos[nn][2]); } - n_bg = 1.68e-4; - n_cl = 5.4e-2; + n_bg = 1e-2; + n_cl = 10; rho_bg = n_bg * mu * MP / DENSITY_UNIT; rho_cl = n_cl * mu * MP / DENSITY_UNIT; - vx_bg = 0.0; + vx_bg = 1000*TIME_UNIT/KPC; // vx_c = -200*TIME_UNIT/KPC; // convert from km/s to kpc/kyr vx_cl = 0.0; vy_bg = vy_cl = 0.0; vz_bg = vz_cl = 0.0; - T_bg = 3e6; + T_bg = 3e7; T_cl = 1e4; p_bg = n_bg * KB * T_bg / PRESSURE_UNIT; p_cl = p_bg; @@ -1414,10 +1414,11 @@ void Grid3D::Clouds() #ifdef DE C.GasEnergy[id] = p_cl / (gama - 1.0); #endif // DE - +#ifdef SCALAR #ifdef DUST C.host[id + H.n_cells * grid_enum::dust_density] = rho_cl * 1e-2; #endif // DUST +#endif } } } diff --git a/src/integrators/VL_3D_cuda.cu b/src/integrators/VL_3D_cuda.cu index bb14be755..a600ecd73 100644 --- a/src/integrators/VL_3D_cuda.cu +++ b/src/integrators/VL_3D_cuda.cu @@ -332,9 +332,11 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int #endif // TEMPERATURE_FLOOR #ifdef SCALAR_FLOOR + #ifdef DUST hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, - grid_enum::dust_density, 1e-5); + grid_enum::dust_density, 1e-10); CudaCheckError(); + #endif #endif // SCALAR_FLOOR return; diff --git a/src/integrators/simple_3D_cuda.cu b/src/integrators/simple_3D_cuda.cu index dd06a2ae6..7a64de081 100644 --- a/src/integrators/simple_3D_cuda.cu +++ b/src/integrators/simple_3D_cuda.cu @@ -187,9 +187,11 @@ void Simple_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, #endif // TEMPERATURE_FLOOR #ifdef SCALAR_FLOOR + #ifdef DUST hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, grid_enum::dust_density, 1e-5); CudaCheckError(); + #endif DUST #endif // SCALAR_FLOOR return; From 098fe222534dcb22a43b0e7b3666a7d6ba9b6161 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Fri, 8 Dec 2023 15:06:08 -0500 Subject: [PATCH 07/40] update initial conditions --- src/grid/cuda_boundaries.cu | 4 ++-- src/grid/initial_conditions.cpp | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/src/grid/cuda_boundaries.cu b/src/grid/cuda_boundaries.cu index 733fb1a85..8e04527ab 100644 --- a/src/grid/cuda_boundaries.cu +++ b/src/grid/cuda_boundaries.cu @@ -312,13 +312,13 @@ __global__ void Wind_Boundary_kernel(Real *c_device, int nx, int ny, int nz, int Real vx, vy, vz, d_0, P_0; n_0 = 1e-2; // same value as n_bg in cloud initial condition function (cm^-3) - T_0 = 3e7; // same value as T_bg in cloud initial condition function (K) + T_0 = 3e6; // same value as T_bg in cloud initial condition function (K) // same values as rho_bg and p_bg in cloud initial condition function d_0 = n_0 * mu * MP / DENSITY_UNIT; P_0 = n_0 * KB * T_0 / PRESSURE_UNIT; - vx = 500 * TIME_UNIT / KPC; // km/s * (cholla unit conversion) + vx = 1000 * TIME_UNIT / KPC; // km/s * (cholla unit conversion) vy = 0.0; vz = 0.0; diff --git a/src/grid/initial_conditions.cpp b/src/grid/initial_conditions.cpp index a1983fb48..0e2ee15ca 100644 --- a/src/grid/initial_conditions.cpp +++ b/src/grid/initial_conditions.cpp @@ -1354,7 +1354,7 @@ void Grid3D::Clouds() vx_cl = 0.0; vy_bg = vy_cl = 0.0; vz_bg = vz_cl = 0.0; - T_bg = 3e7; + T_bg = 3e6; T_cl = 1e4; p_bg = n_bg * KB * T_bg / PRESSURE_UNIT; p_cl = p_bg; From 9b7e6147cd46926c2dfed9a23ebc5a642e33c76d Mon Sep 17 00:00:00 2001 From: helenarichie Date: Thu, 14 Dec 2023 08:24:31 -0500 Subject: [PATCH 08/40] save simulation setup --- src/grid/cuda_boundaries.cu | 2 +- src/grid/initial_conditions.cpp | 7 ++++--- 2 files changed, 5 insertions(+), 4 deletions(-) diff --git a/src/grid/cuda_boundaries.cu b/src/grid/cuda_boundaries.cu index 8e04527ab..9ad4a80e6 100644 --- a/src/grid/cuda_boundaries.cu +++ b/src/grid/cuda_boundaries.cu @@ -312,7 +312,7 @@ __global__ void Wind_Boundary_kernel(Real *c_device, int nx, int ny, int nz, int Real vx, vy, vz, d_0, P_0; n_0 = 1e-2; // same value as n_bg in cloud initial condition function (cm^-3) - T_0 = 3e6; // same value as T_bg in cloud initial condition function (K) + T_0 = 3e7; // same value as T_bg in cloud initial condition function (K) // same values as rho_bg and p_bg in cloud initial condition function d_0 = n_0 * mu * MP / DENSITY_UNIT; diff --git a/src/grid/initial_conditions.cpp b/src/grid/initial_conditions.cpp index 0e2ee15ca..c25f24c31 100644 --- a/src/grid/initial_conditions.cpp +++ b/src/grid/initial_conditions.cpp @@ -1346,7 +1346,7 @@ void Grid3D::Clouds() } n_bg = 1e-2; - n_cl = 10; + n_cl = 50; rho_bg = n_bg * mu * MP / DENSITY_UNIT; rho_cl = n_cl * mu * MP / DENSITY_UNIT; vx_bg = 1000*TIME_UNIT/KPC; @@ -1354,10 +1354,11 @@ void Grid3D::Clouds() vx_cl = 0.0; vy_bg = vy_cl = 0.0; vz_bg = vz_cl = 0.0; - T_bg = 3e6; + T_bg = 3e7; T_cl = 1e4; p_bg = n_bg * KB * T_bg / PRESSURE_UNIT; - p_cl = p_bg; + // p_cl = p_bg; + p_cl = n_cl * KB * T_cl / PRESSURE_UNIT; istart = H.n_ghost; iend = H.nx - H.n_ghost; From 4077eddf208d335ca7f4f8710fb0d97a867cd763 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Thu, 14 Dec 2023 08:31:45 -0500 Subject: [PATCH 09/40] resolve merge conflict --- src/integrators/VL_3D_cuda.cu | 9 ++------- 1 file changed, 2 insertions(+), 7 deletions(-) diff --git a/src/integrators/VL_3D_cuda.cu b/src/integrators/VL_3D_cuda.cu index 7d10dff5a..d562707ea 100644 --- a/src/integrators/VL_3D_cuda.cu +++ b/src/integrators/VL_3D_cuda.cu @@ -194,19 +194,14 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int // Step 3: Update the conserved variables half a timestep hipLaunchKernelGGL(Update_Conserved_Variables_3D_half, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, dev_conserved_half, -<<<<<<< HEAD - F_x, F_y, F_z, nx, ny, nz, n_ghost, dx, dy, dz, 0.5 * dt, gama, n_fields); - CudaCheckError(); + F_x, F_y, F_z, nx, ny, nz, n_ghost, dx, dy, dz, 0.5 * dt, gama, n_fields, density_floor); + GPU_Error_Check(); #ifdef DENSITY_FLOOR hipLaunchKernelGGL(Apply_Density_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved_half, nx, ny, nz, n_ghost, density_floor); #endif // DENSITY_FLOOR -======= - F_x, F_y, F_z, nx, ny, nz, n_ghost, dx, dy, dz, 0.5 * dt, gama, n_fields, density_floor); - GPU_Error_Check(); ->>>>>>> dev #ifdef MHD // Update the magnetic fields hipLaunchKernelGGL(mhd::Update_Magnetic_Field_3D, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, dev_conserved_half, From 3cd6b2ef9e04369686229035727e3d93aff1ee14 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Fri, 15 Dec 2023 08:54:31 -0500 Subject: [PATCH 10/40] update from cell averaging PR --- builds/make.type.dust | 4 ++-- src/integrators/VL_3D_cuda.cu | 7 ++++--- 2 files changed, 6 insertions(+), 5 deletions(-) diff --git a/builds/make.type.dust b/builds/make.type.dust index b9813674c..45a11f61b 100644 --- a/builds/make.type.dust +++ b/builds/make.type.dust @@ -14,8 +14,8 @@ DFLAGS += -DHLLC # DFLAGS += -DDE DFLAGS += -DAVERAGE_SLOW_CELLS -DFLAGS += -DTEMPERATURE_FLOOR -DFLAGS += -DSCALAR_FLOOR +# DFLAGS += -DTEMPERATURE_FLOOR +# DFLAGS += -DSCALAR_FLOOR ifeq ($(findstring cosmology,$(TYPE)),cosmology) DFLAGS += -DSIMPLE diff --git a/src/integrators/VL_3D_cuda.cu b/src/integrators/VL_3D_cuda.cu index d562707ea..f299eda3c 100644 --- a/src/integrators/VL_3D_cuda.cu +++ b/src/integrators/VL_3D_cuda.cu @@ -32,7 +32,7 @@ __global__ void Update_Conserved_Variables_3D_half(Real *dev_conserved, Real *dev_conserved_half, Real *dev_F_x, Real *dev_F_y, Real *dev_F_z, int nx, int ny, int nz, int n_ghost, - Real dx, Real dy, Real dz, Real dt, Real gamma, int n_fields); + Real dx, Real dy, Real dz, Real dt, Real gamma, int n_fields, Real density_floor); 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, @@ -335,7 +335,7 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int #ifdef DUST hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, grid_enum::dust_density, 1e-10); - CudaCheckError(); + GPU_Error_Check(); #endif #endif // SCALAR_FLOOR @@ -361,7 +361,8 @@ void Free_Memory_VL_3D() __global__ void Update_Conserved_Variables_3D_half(Real *dev_conserved, Real *dev_conserved_half, Real *dev_F_x, Real *dev_F_y, Real *dev_F_z, int nx, int ny, int nz, int n_ghost, - Real dx, Real dy, Real dz, Real dt, Real gamma, int n_fields) + Real dx, Real dy, Real dz, Real dt, Real gamma, int n_fields, + Real density_floor) { Real dtodx = dt / dx; Real dtody = dt / dy; From c1ccee3ca8783a5b4aab6a8dd090a7e07c040293 Mon Sep 17 00:00:00 2001 From: Helena Richie Date: Wed, 3 Jan 2024 12:30:59 -0500 Subject: [PATCH 11/40] update branch --- src/grid/initial_conditions.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/src/grid/initial_conditions.cpp b/src/grid/initial_conditions.cpp index 432321113..c9d22727a 100644 --- a/src/grid/initial_conditions.cpp +++ b/src/grid/initial_conditions.cpp @@ -1346,7 +1346,7 @@ void Grid3D::Clouds() } n_bg = 1e-2; - n_cl = 50; + n_cl = 10; rho_bg = n_bg * mu * MP / DENSITY_UNIT; rho_cl = n_cl * mu * MP / DENSITY_UNIT; vx_bg = 1000*TIME_UNIT/KPC; @@ -1357,8 +1357,8 @@ void Grid3D::Clouds() T_bg = 3e7; T_cl = 1e4; p_bg = n_bg * KB * T_bg / PRESSURE_UNIT; - // p_cl = p_bg; - p_cl = n_cl * KB * T_cl / PRESSURE_UNIT; + p_cl = p_bg; + // p_cl = n_cl * KB * T_cl / PRESSURE_UNIT; istart = H.n_ghost; iend = H.nx - H.n_ghost; From 5b8525187f8a5dc965d4a999e4285f57ca8a7f5d Mon Sep 17 00:00:00 2001 From: Helena Richie Date: Thu, 4 Jan 2024 08:28:49 -0500 Subject: [PATCH 12/40] update build --- builds/make.type.dust | 11 +++-------- src/dust/dust_cuda.cu | 2 +- 2 files changed, 4 insertions(+), 9 deletions(-) diff --git a/builds/make.type.dust b/builds/make.type.dust index 45a11f61b..a95912ea0 100644 --- a/builds/make.type.dust +++ b/builds/make.type.dust @@ -9,20 +9,15 @@ MPI_GPU ?= DFLAGS += -DCUDA DFLAGS += -DMPI_CHOLLA DFLAGS += -DPRECISION=2 -DFLAGS += -DPPMP +DFLAGS += -DPLMP DFLAGS += -DHLLC -# DFLAGS += -DDE +DFLAGS += -DDE DFLAGS += -DAVERAGE_SLOW_CELLS # DFLAGS += -DTEMPERATURE_FLOOR -# DFLAGS += -DSCALAR_FLOOR +DFLAGS += -DSCALAR_FLOOR -ifeq ($(findstring cosmology,$(TYPE)),cosmology) -DFLAGS += -DSIMPLE -else DFLAGS += -DVL -# DFLAGS += -DSIMPLE -endif # Evolve additional scalars DFLAGS += -DSCALAR diff --git a/src/dust/dust_cuda.cu b/src/dust/dust_cuda.cu index 7ca48a9fd..d4c6191ec 100644 --- a/src/dust/dust_cuda.cu +++ b/src/dust/dust_cuda.cu @@ -83,7 +83,7 @@ __global__ void Dust_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_g velocity_z = dev_conserved[id + n_cells * grid_enum::momentum_z] / density_gas; #ifdef DE energy_gas = dev_conserved[id + n_cells * grid_enum::GasEnergy] / density_gas; - energy_gas = fmax(ge, (Real)TINY_NUMBER); + energy_gas = fmax(energy_gas, (Real)TINY_NUMBER); #endif // DE // calculate physical quantities From f5dc360bdf2c2a5e649b665af7add87ea2d6aa1f Mon Sep 17 00:00:00 2001 From: helenarichie Date: Thu, 4 Jan 2024 09:00:07 -0500 Subject: [PATCH 13/40] hard-code cloud temp --- src/grid/cuda_boundaries.cu | 2 +- src/grid/initial_conditions.cpp | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/src/grid/cuda_boundaries.cu b/src/grid/cuda_boundaries.cu index 675f14ae5..715840322 100644 --- a/src/grid/cuda_boundaries.cu +++ b/src/grid/cuda_boundaries.cu @@ -312,7 +312,7 @@ __global__ void Wind_Boundary_kernel(Real *c_device, int nx, int ny, int nz, int Real vx, vy, vz, d_0, P_0; n_0 = 1e-2; // same value as n_bg in cloud initial condition function (cm^-3) - T_0 = 3e7; // same value as T_bg in cloud initial condition function (K) + T_0 = 3e6; // same value as T_bg in cloud initial condition function (K) // same values as rho_bg and p_bg in cloud initial condition function d_0 = n_0 * mu * MP / DENSITY_UNIT; diff --git a/src/grid/initial_conditions.cpp b/src/grid/initial_conditions.cpp index 432321113..658362954 100644 --- a/src/grid/initial_conditions.cpp +++ b/src/grid/initial_conditions.cpp @@ -1346,7 +1346,7 @@ void Grid3D::Clouds() } n_bg = 1e-2; - n_cl = 50; + n_cl = 10; rho_bg = n_bg * mu * MP / DENSITY_UNIT; rho_cl = n_cl * mu * MP / DENSITY_UNIT; vx_bg = 1000*TIME_UNIT/KPC; @@ -1354,7 +1354,7 @@ void Grid3D::Clouds() vx_cl = 0.0; vy_bg = vy_cl = 0.0; vz_bg = vz_cl = 0.0; - T_bg = 3e7; + T_bg = 3e6; T_cl = 1e4; p_bg = n_bg * KB * T_bg / PRESSURE_UNIT; // p_cl = p_bg; From e12b846a388f3429920196f1088773babb540c82 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Thu, 4 Jan 2024 14:01:33 -0500 Subject: [PATCH 14/40] update build --- src/grid/initial_conditions.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/grid/initial_conditions.cpp b/src/grid/initial_conditions.cpp index 6389b8bf7..658362954 100644 --- a/src/grid/initial_conditions.cpp +++ b/src/grid/initial_conditions.cpp @@ -1357,8 +1357,8 @@ void Grid3D::Clouds() T_bg = 3e6; T_cl = 1e4; p_bg = n_bg * KB * T_bg / PRESSURE_UNIT; - p_cl = p_bg; - // p_cl = n_cl * KB * T_cl / PRESSURE_UNIT; + // p_cl = p_bg; + p_cl = n_cl * KB * T_cl / PRESSURE_UNIT; istart = H.n_ghost; iend = H.nx - H.n_ghost; From 3fec9785d8bfd6654ad9045cfc8135bd9822a218 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Sat, 6 Jan 2024 16:37:34 -0500 Subject: [PATCH 15/40] update build --- builds/make.type.dust | 4 +--- src/global/global.h | 2 +- src/grid/grid3D.cpp | 2 +- src/hydro/hydro_cuda.cu | 3 ++- 4 files changed, 5 insertions(+), 6 deletions(-) diff --git a/builds/make.type.dust b/builds/make.type.dust index a95912ea0..35d595a37 100644 --- a/builds/make.type.dust +++ b/builds/make.type.dust @@ -14,7 +14,7 @@ DFLAGS += -DHLLC DFLAGS += -DDE DFLAGS += -DAVERAGE_SLOW_CELLS -# DFLAGS += -DTEMPERATURE_FLOOR +DFLAGS += -DTEMPERATURE_FLOOR DFLAGS += -DSCALAR_FLOOR DFLAGS += -DVL @@ -25,8 +25,6 @@ DFLAGS += -DSCALAR # Define dust macro DFLAGS += -DDUST -DFLAGS += -DSCALAR_FLOOR - # Apply the cooling in the GPU from precomputed tables DFLAGS += -DCOOLING_GPU DFLAGS += -DCLOUDY_COOLING diff --git a/src/global/global.h b/src/global/global.h index dbb976d8c..18b08b71f 100644 --- a/src/global/global.h +++ b/src/global/global.h @@ -50,7 +50,7 @@ typedef double Real; #define LOG_FILE_NAME "run_output.log" // Conserved Floor Values -#define TEMP_FLOOR 10 +#define TEMP_FLOOR 100 #define DENS_FLOOR 1e-5 // in code units // Parameters for Enzo dual Energy Condition diff --git a/src/grid/grid3D.cpp b/src/grid/grid3D.cpp index a76417321..4638942eb 100644 --- a/src/grid/grid3D.cpp +++ b/src/grid/grid3D.cpp @@ -157,7 +157,7 @@ void Grid3D::Initialize(struct Parameters *P) C_cfl = 0.3; #ifdef AVERAGE_SLOW_CELLS - H.min_dt_slow = 1e-5; // Initialize the minumum dt to a tiny number + H.min_dt_slow = 0.024; // Initialize the minumum dt to a tiny number #endif // AVERAGE_SLOW_CELLS #ifndef MPI_CHOLLA diff --git a/src/hydro/hydro_cuda.cu b/src/hydro/hydro_cuda.cu index bff6df5ba..e50cdd3f0 100644 --- a/src/hydro/hydro_cuda.cu +++ b/src/hydro/hydro_cuda.cu @@ -379,7 +379,8 @@ __global__ void Update_Conserved_Variables_3D(Real *dev_conserved, Real *Q_Lx, R #endif // GRAVITY - #if !(defined(DENSITY_FLOOR) && defined(TEMPERATURE_FLOOR)) + // #if !(defined(DENSITY_FLOOR) && defined(TEMPERATURE_FLOOR)) + #if !(defined(DENSITY_FLOOR)) 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 %e %e %e\n", xid + x_off, yid + y_off, zid + z_off, From 328770bda11c5d3d1b3253bf3c14f9c94fa39e16 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Sat, 13 Jan 2024 15:15:30 -0500 Subject: [PATCH 16/40] clean up initial conditions --- src/grid/initial_conditions.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/grid/initial_conditions.cpp b/src/grid/initial_conditions.cpp index 658362954..87c5d392f 100644 --- a/src/grid/initial_conditions.cpp +++ b/src/grid/initial_conditions.cpp @@ -1419,7 +1419,7 @@ void Grid3D::Clouds() #ifdef DUST C.host[id + H.n_cells * grid_enum::dust_density] = rho_cl * 1e-2; #endif // DUST -#endif +#endif // SCAlAR } } } From 2f3ef5fed87eeba6e9c25b9424fca7e15cacf829 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Sat, 13 Jan 2024 15:23:16 -0500 Subject: [PATCH 17/40] reset to default settings --- builds/make.type.dust | 18 +++++++++--------- src/global/global.h | 2 +- src/grid/cuda_boundaries.cu | 2 +- src/grid/grid3D.cpp | 2 +- 4 files changed, 12 insertions(+), 12 deletions(-) diff --git a/builds/make.type.dust b/builds/make.type.dust index 35d595a37..a6cf1bdd0 100644 --- a/builds/make.type.dust +++ b/builds/make.type.dust @@ -15,25 +15,25 @@ DFLAGS += -DHLLC DFLAGS += -DDE DFLAGS += -DAVERAGE_SLOW_CELLS DFLAGS += -DTEMPERATURE_FLOOR -DFLAGS += -DSCALAR_FLOOR DFLAGS += -DVL # Evolve additional scalars -DFLAGS += -DSCALAR +DFLAGS += -DSCALAR +DFLAGS += -DSCALAR_FLOOR # Define dust macro -DFLAGS += -DDUST +DFLAGS += -DDUST # Apply the cooling in the GPU from precomputed tables -DFLAGS += -DCOOLING_GPU -DFLAGS += -DCLOUDY_COOLING +DFLAGS += -DCOOLING_GPU +DFLAGS += -DCLOUDY_COOLING #Measure the Timing of the different stages -#DFLAGS += -DCPU_TIME +#DFLAGS += -DCPU_TIME -DFLAGS += -DSLICES -DFLAGS += -DPROJECTION +DFLAGS += -DSLICES +DFLAGS += -DPROJECTION DFLAGS += $(OUTPUT) @@ -41,4 +41,4 @@ DFLAGS += $(OUTPUT) #and the MPI transfers are done from the GPU #If not specified, MPI_GPU is off by default #This is set in the system make.host file -DFLAGS += $(MPI_GPU) +DFLAGS += $(MPI_GPU) \ No newline at end of file diff --git a/src/global/global.h b/src/global/global.h index 18b08b71f..f7030b563 100644 --- a/src/global/global.h +++ b/src/global/global.h @@ -50,7 +50,7 @@ typedef double Real; #define LOG_FILE_NAME "run_output.log" // Conserved Floor Values -#define TEMP_FLOOR 100 +#define TEMP_FLOOR 1e-3 #define DENS_FLOOR 1e-5 // in code units // Parameters for Enzo dual Energy Condition diff --git a/src/grid/cuda_boundaries.cu b/src/grid/cuda_boundaries.cu index 715840322..baf846d3c 100644 --- a/src/grid/cuda_boundaries.cu +++ b/src/grid/cuda_boundaries.cu @@ -318,7 +318,7 @@ __global__ void Wind_Boundary_kernel(Real *c_device, int nx, int ny, int nz, int d_0 = n_0 * mu * MP / DENSITY_UNIT; P_0 = n_0 * KB * T_0 / PRESSURE_UNIT; - vx = 1000 * TIME_UNIT / KPC; // km/s * (cholla unit conversion) + vx = 100 * TIME_UNIT / KPC; // km/s * (cholla unit conversion) vy = 0.0; vz = 0.0; diff --git a/src/grid/grid3D.cpp b/src/grid/grid3D.cpp index 4638942eb..a76417321 100644 --- a/src/grid/grid3D.cpp +++ b/src/grid/grid3D.cpp @@ -157,7 +157,7 @@ void Grid3D::Initialize(struct Parameters *P) C_cfl = 0.3; #ifdef AVERAGE_SLOW_CELLS - H.min_dt_slow = 0.024; // Initialize the minumum dt to a tiny number + H.min_dt_slow = 1e-5; // Initialize the minumum dt to a tiny number #endif // AVERAGE_SLOW_CELLS #ifndef MPI_CHOLLA From 094f1e9ad96c0f04b60aa9a782c4de11356c2e06 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Sat, 13 Jan 2024 17:24:46 -0500 Subject: [PATCH 18/40] get floor values from paramter file and remove density floor kernel --- builds/make.type.dust | 3 +++ src/global/global.cpp | 12 +++++++++++ src/global/global.h | 13 ++++++++---- src/grid/grid3D.cpp | 18 ++++++++++------ src/grid/grid3D.h | 3 ++- src/grid/initial_conditions.cpp | 15 +++++++------ src/hydro/hydro_cuda.cu | 35 ++----------------------------- src/hydro/hydro_cuda.h | 2 -- src/hydro/hydro_cuda_tests.cu | 1 - src/integrators/VL_3D_cuda.cu | 9 ++------ src/integrators/VL_3D_cuda.h | 2 +- src/integrators/simple_3D_cuda.cu | 4 ++-- src/integrators/simple_3D_cuda.h | 2 +- 13 files changed, 53 insertions(+), 66 deletions(-) diff --git a/builds/make.type.dust b/builds/make.type.dust index a6cf1bdd0..24a765302 100644 --- a/builds/make.type.dust +++ b/builds/make.type.dust @@ -15,6 +15,7 @@ DFLAGS += -DHLLC DFLAGS += -DDE DFLAGS += -DAVERAGE_SLOW_CELLS DFLAGS += -DTEMPERATURE_FLOOR +DFLAGS += -DDENSITY_FLOOR DFLAGS += -DVL @@ -37,6 +38,8 @@ DFLAGS += -DPROJECTION DFLAGS += $(OUTPUT) +DFLAGS += -DOUTPUT_ALWAYS + #Select if the Hydro Conserved data will reside in the GPU #and the MPI transfers are done from the GPU #If not specified, MPI_GPU is off by default diff --git a/src/global/global.cpp b/src/global/global.cpp index 89347cbda..d0f898739 100644 --- a/src/global/global.cpp +++ b/src/global/global.cpp @@ -428,6 +428,18 @@ void Parse_Param(char *name, char *value, struct Parameters *parms) } else if (strcmp(name, "UVB_rates_file") == 0) { strncpy(parms->UVB_rates_file, value, MAXLEN); #endif +# ifdef TEMPERATURE_FLOOR + } else if (strcmp(name, "temperature_floor") == 0) { + parms->temperature_floor = atof(value); +# endif +# ifdef DENSITY_FLOOR + } else if (strcmp(name, "density_floor") == 0) { + parms->density_floor = atof(value); +# endif +# ifdef SCALAR_FLOOR + } else if (strcmp(name, "scalar_floor") == 0) { + parms->scalar_floor = atof(value); +# endif #ifdef ANALYSIS } else if (strcmp(name, "analysis_scale_outputs_file") == 0) { strncpy(parms->analysis_scale_outputs_file, value, MAXLEN); diff --git a/src/global/global.h b/src/global/global.h index f7030b563..6c9524001 100644 --- a/src/global/global.h +++ b/src/global/global.h @@ -49,10 +49,6 @@ typedef double Real; #define LOG_FILE_NAME "run_output.log" -// Conserved Floor Values -#define TEMP_FLOOR 1e-3 -#define DENS_FLOOR 1e-5 // in code units - // Parameters for Enzo dual Energy Condition // - Prior to GH PR #356, DE_ETA_1 nominally had a value of 0.001 in all // simulations (in practice, the value of DE_ETA_1 had minimal significance @@ -325,6 +321,15 @@ struct Parameters { char UVB_rates_file[MAXLEN]; // File for the UVB photoheating and // photoionization rates of HI, HeI and HeII #endif +#ifdef TEMPERATURE_FLOOR + Real temperature_floor; +#endif +#ifdef DENSITY_FLOOR + Real density_floor; +#endif +#ifdef SCALAR_FLOOR + Real scalar_floor; +#endif #ifdef ANALYSIS char analysis_scale_outputs_file[MAXLEN]; // File for the scale_factor output // values for cosmological diff --git a/src/grid/grid3D.cpp b/src/grid/grid3D.cpp index a76417321..d844784d5 100644 --- a/src/grid/grid3D.cpp +++ b/src/grid/grid3D.cpp @@ -258,16 +258,22 @@ void Grid3D::Initialize(struct Parameters *P) #endif /*ROTATED_PROJECTION*/ // Values for lower limit for density and temperature +#ifdef TEMPERATURE_FLOOR + H.temperature_floor = P->temperature_floor; +#else + H.temperature_floor = 0.0; +#endif + #ifdef DENSITY_FLOOR - H.density_floor = DENS_FLOOR; + H.density_floor = P->density_floor; #else H.density_floor = 0.0; #endif -#ifdef TEMPERATURE_FLOOR - H.temperature_floor = TEMP_FLOOR; +#ifdef SCALAR_FLOOR + H.scalar_floor = P->scalar_floor; #else - H.temperature_floor = 0.0; + H.scalar_floor = 0.0; #endif #ifdef COSMOLOGY @@ -461,12 +467,12 @@ 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, density_floor, U_floor, - C.Grav_potential); + C.Grav_potential, H.scalar_floor); #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, H.dz, H.xbound, H.ybound, H.zbound, H.dt, H.n_fields, H.custom_grav, density_floor, - U_floor, C.Grav_potential); + U_floor, C.Grav_potential, H.scalar_floor); #endif // SIMPLE #endif } else { diff --git a/src/grid/grid3D.h b/src/grid/grid3D.h index aff94c898..5354acf0a 100644 --- a/src/grid/grid3D.h +++ b/src/grid/grid3D.h @@ -231,8 +231,9 @@ struct Header { int custom_grav; // Values for lower limit for density and temperature - Real density_floor; Real temperature_floor; + Real density_floor; + Real scalar_floor; Real Ekin_avrg; diff --git a/src/grid/initial_conditions.cpp b/src/grid/initial_conditions.cpp index 87c5d392f..0e5aabc90 100644 --- a/src/grid/initial_conditions.cpp +++ b/src/grid/initial_conditions.cpp @@ -1324,7 +1324,7 @@ void Grid3D::Clouds() Real p_bg, p_cl; // background and cloud pressure Real mu = 0.6; // mean atomic weight int N_cl = 1; // number of clouds - Real R_cl = .1; // cloud radius in code units (kpc) + Real R_cl = 2.5; // cloud radius in code units (kpc) Real cl_pos[N_cl][3]; // array of cloud positions Real r; @@ -1339,17 +1339,17 @@ void Grid3D::Clouds() // single centered cloud setup for (int nn = 0; nn < N_cl; nn++) { - cl_pos[nn][0] = 0.075 * H.xdglobal; + cl_pos[nn][0] = 0.5 * H.xdglobal; cl_pos[nn][1] = 0.5 * H.ydglobal; cl_pos[nn][2] = 0.5 * H.zdglobal; printf("Cloud positions: %f %f %f\n", cl_pos[nn][0], cl_pos[nn][1], cl_pos[nn][2]); } - n_bg = 1e-2; - n_cl = 10; + n_bg = 1.68e-4; + n_cl = 5.4e-2; rho_bg = n_bg * mu * MP / DENSITY_UNIT; rho_cl = n_cl * mu * MP / DENSITY_UNIT; - vx_bg = 1000*TIME_UNIT/KPC; + vx_bg = 0.0; // vx_c = -200*TIME_UNIT/KPC; // convert from km/s to kpc/kyr vx_cl = 0.0; vy_bg = vy_cl = 0.0; @@ -1357,8 +1357,7 @@ void Grid3D::Clouds() T_bg = 3e6; T_cl = 1e4; p_bg = n_bg * KB * T_bg / PRESSURE_UNIT; - // p_cl = p_bg; - p_cl = n_cl * KB * T_cl / PRESSURE_UNIT; + p_cl = p_bg; istart = H.n_ghost; iend = H.nx - H.n_ghost; @@ -1419,7 +1418,7 @@ void Grid3D::Clouds() #ifdef DUST C.host[id + H.n_cells * grid_enum::dust_density] = rho_cl * 1e-2; #endif // DUST -#endif // SCAlAR +#endif // SCALAR } } } diff --git a/src/hydro/hydro_cuda.cu b/src/hydro/hydro_cuda.cu index e50cdd3f0..125e851b4 100644 --- a/src/hydro/hydro_cuda.cu +++ b/src/hydro/hydro_cuda.cu @@ -379,8 +379,7 @@ __global__ void Update_Conserved_Variables_3D(Real *dev_conserved, Real *Q_Lx, R #endif // GRAVITY - // #if !(defined(DENSITY_FLOOR) && defined(TEMPERATURE_FLOOR)) - #if !(defined(DENSITY_FLOOR)) + #if !(defined(DENSITY_FLOOR) && defined(TEMPERATURE_FLOOR)) 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 %e %e %e\n", xid + x_off, yid + y_off, zid + z_off, @@ -1140,37 +1139,6 @@ __global__ void Apply_Temperature_Floor(Real *dev_conserved, int nx, int ny, int } } -__global__ void Apply_Density_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, Real density_floor) -{ - int id, xid, yid, zid, n_cells; - Real density_init; // variable to store the value of the scalar before a floor is applied - n_cells = nx * ny * nz; - - // get a global thread ID - id = threadIdx.x + blockIdx.x * blockDim.x; - zid = id / (nx * ny); - yid = (id - zid * nx * ny) / nx; - xid = id - zid * nx * ny - yid * nx; - - // threads corresponding to real cells do the calculation - 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) { - density_init = dev_conserved[id + n_cells * grid_enum::density]; - - if (density_init < density_floor) { - printf("###Thread density change %f -> %f \n", density_init, density_floor); - dev_conserved[id] = density_floor; - // Scale the conserved values to the new density - dev_conserved[id + n_cells * grid_enum::momentum_x] *= (density_floor / density_init); - dev_conserved[id + n_cells * grid_enum::momentum_y] *= (density_floor / density_init); - dev_conserved[id + n_cells * grid_enum::momentum_z] *= (density_floor / density_init); - dev_conserved[id + n_cells * grid_enum::Energy] *= (density_floor / density_init); - #ifdef DE - dev_conserved[id + n_cells * grid_enum::GasEnergy] *= (density_floor / density_init); - #endif // DE - } - } -} __device__ Real Average_Cell_Single_Field(int field_indx, int i, int j, int k, int nx, int ny, int nz, int ncells, Real *conserved) { @@ -1305,6 +1273,7 @@ __global__ void Apply_Scalar_Floor(Real *dev_conserved, int nx, int ny, int nz, scalar = dev_conserved[id + n_cells * field_num]; if (scalar < scalar_floor) { + printf("###Thread scalar change %f -> %f \n", scalar, scalar_floor); dev_conserved[id + n_cells * field_num] = scalar_floor; } } diff --git a/src/hydro/hydro_cuda.h b/src/hydro/hydro_cuda.h index 92a997790..8fcfbba05 100644 --- a/src/hydro/hydro_cuda.h +++ b/src/hydro/hydro_cuda.h @@ -88,8 +88,6 @@ __global__ void Average_Slow_Cells_3D(Real *dev_conserved, int nx, int ny, int n __global__ void Apply_Temperature_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real U_floor); -__global__ void Apply_Density_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, Real density_floor); - __global__ void Apply_Scalar_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, Real scalar_floor); diff --git a/src/hydro/hydro_cuda_tests.cu b/src/hydro/hydro_cuda_tests.cu index 0403478e5..0796a3064 100644 --- a/src/hydro/hydro_cuda_tests.cu +++ b/src/hydro/hydro_cuda_tests.cu @@ -148,7 +148,6 @@ TEST(tMHDMhdInverseCrossingTime, CorrectInputExpectCorrectOutput) TEST(tHYDROScalarFloor, CorrectInputExpectCorrectOutput) { - // Call the function we are testing int num_blocks = 1; dim3 dim1dGrid(num_blocks, 1, 1); dim3 dim1dBlock(TPB, 1, 1); diff --git a/src/integrators/VL_3D_cuda.cu b/src/integrators/VL_3D_cuda.cu index f299eda3c..25462fa91 100644 --- a/src/integrators/VL_3D_cuda.cu +++ b/src/integrators/VL_3D_cuda.cu @@ -37,7 +37,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 U_floor, - Real *host_grav_potential) + Real *host_grav_potential, Real scalar_floor) { // Here, *dev_conserved contains the entire // set of conserved variables on the grid @@ -197,11 +197,6 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int F_x, F_y, F_z, nx, ny, nz, n_ghost, dx, dy, dz, 0.5 * dt, gama, n_fields, density_floor); GPU_Error_Check(); - #ifdef DENSITY_FLOOR - hipLaunchKernelGGL(Apply_Density_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved_half, nx, ny, nz, n_ghost, - density_floor); - #endif // DENSITY_FLOOR - #ifdef MHD // Update the magnetic fields hipLaunchKernelGGL(mhd::Update_Magnetic_Field_3D, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, dev_conserved_half, @@ -334,7 +329,7 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int #ifdef SCALAR_FLOOR #ifdef DUST hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, - grid_enum::dust_density, 1e-10); + grid_enum::dust_density, scalar_floor); GPU_Error_Check(); #endif #endif // SCALAR_FLOOR diff --git a/src/integrators/VL_3D_cuda.h b/src/integrators/VL_3D_cuda.h index 3f2cf8d75..4104493bc 100644 --- a/src/integrators/VL_3D_cuda.h +++ b/src/integrators/VL_3D_cuda.h @@ -11,7 +11,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 U_floor, - Real *host_grav_potential); + Real *host_grav_potential, Real scalar_floor); void Free_Memory_VL_3D(); diff --git a/src/integrators/simple_3D_cuda.cu b/src/integrators/simple_3D_cuda.cu index c572dcd97..d3beb14b2 100644 --- a/src/integrators/simple_3D_cuda.cu +++ b/src/integrators/simple_3D_cuda.cu @@ -27,7 +27,7 @@ void Simple_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 U_floor, - Real *host_grav_potential) + Real *host_grav_potential, Real scalar_floor) { // Here, *dev_conserved contains the entire // set of conserved variables on the grid @@ -189,7 +189,7 @@ void Simple_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, #ifdef SCALAR_FLOOR #ifdef DUST hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, - grid_enum::dust_density, 1e-5); + grid_enum::dust_density, scalar_floor); CudaCheckError(); #endif DUST #endif // SCALAR_FLOOR diff --git a/src/integrators/simple_3D_cuda.h b/src/integrators/simple_3D_cuda.h index 585c553ba..e2cea247e 100644 --- a/src/integrators/simple_3D_cuda.h +++ b/src/integrators/simple_3D_cuda.h @@ -12,7 +12,7 @@ void Simple_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 U_floor, - Real *host_grav_potential); + Real *host_grav_potential, Real scalar_floor); void Free_Memory_Simple_3D(); From d783bb74f1e8542709cebd8f9b7181b9f2c9afb0 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Sat, 13 Jan 2024 17:26:42 -0500 Subject: [PATCH 19/40] run clang format --- src/analysis/feedback_analysis.cpp | 2 +- src/chemistry_gpu/chemistry_functions.cpp | 2 +- src/cooling_grackle/cool_grackle.cpp | 2 +- src/global/global.cpp | 12 +++++----- src/gravity/gravity_functions.cpp | 8 +++---- src/gravity/gravity_functions_gpu.cu | 4 ++-- src/grid/grid3D.cpp | 10 ++++----- src/grid/initial_conditions.cpp | 26 +++++++++++----------- src/integrators/VL_3D_cuda.cu | 9 ++++---- src/integrators/simple_3D_cuda.cu | 4 ++-- src/io/io.cpp | 4 ++-- src/particles/io_particles.cpp | 14 ++++++------ src/particles/particles_3D.cpp | 14 ++++++------ src/particles/particles_boundaries_cpu.cpp | 4 ++-- src/reconstruction/plmp_cuda.cu | 2 +- src/reconstruction/ppmc_cuda_tests.cu | 2 +- src/reconstruction/ppmp_cuda.cu | 2 +- src/reconstruction/reconstruction_tests.cu | 8 +++++-- src/system_tests/hydro_system_tests.cpp | 4 ++-- 19 files changed, 69 insertions(+), 64 deletions(-) diff --git a/src/analysis/feedback_analysis.cpp b/src/analysis/feedback_analysis.cpp index 3dab7b6da..5fe0e7543 100644 --- a/src/analysis/feedback_analysis.cpp +++ b/src/analysis/feedback_analysis.cpp @@ -87,7 +87,7 @@ void FeedbackAnalysis::Compute_Gas_Velocity_Dispersion(Grid3D& G) #ifdef MPI_CHOLLA MPI_Allreduce(&partial_mass, &total_mass, 1, MPI_CHREAL, MPI_SUM, world); #else - total_mass = partial_mass; + total_mass = partial_mass; #endif for (k = G.H.n_ghost; k < G.H.nz - G.H.n_ghost; k++) { diff --git a/src/chemistry_gpu/chemistry_functions.cpp b/src/chemistry_gpu/chemistry_functions.cpp index 7999a6d55..65e3af691 100644 --- a/src/chemistry_gpu/chemistry_functions.cpp +++ b/src/chemistry_gpu/chemistry_functions.cpp @@ -228,7 +228,7 @@ void Grid3D::Update_Chemistry() #ifdef COSMOLOGY Chem.H.current_z = Cosmo.current_z; #else - Chem.H.current_z = 0; + Chem.H.current_z = 0; #endif Do_Chemistry_Update(C.device, H.nx, H.ny, H.nz, H.n_ghost, H.n_fields, H.dt, Chem.H); diff --git a/src/cooling_grackle/cool_grackle.cpp b/src/cooling_grackle/cool_grackle.cpp index a7f5c36cb..c5f2a8078 100644 --- a/src/cooling_grackle/cool_grackle.cpp +++ b/src/cooling_grackle/cool_grackle.cpp @@ -89,7 +89,7 @@ void Cool_GK::Initialize(struct Parameters *P, Cosmology &Cosmo) data->metal_cooling = 1; // metal cooling off #else chprintf("WARNING: Metal Cooling is Off. \n"); - data->metal_cooling = 0; // metal cooling off + data->metal_cooling = 0; // metal cooling off #endif #ifdef PARALLEL_OMP diff --git a/src/global/global.cpp b/src/global/global.cpp index d0f898739..6579d9955 100644 --- a/src/global/global.cpp +++ b/src/global/global.cpp @@ -428,18 +428,18 @@ void Parse_Param(char *name, char *value, struct Parameters *parms) } else if (strcmp(name, "UVB_rates_file") == 0) { strncpy(parms->UVB_rates_file, value, MAXLEN); #endif -# ifdef TEMPERATURE_FLOOR +#ifdef TEMPERATURE_FLOOR } else if (strcmp(name, "temperature_floor") == 0) { parms->temperature_floor = atof(value); -# endif -# ifdef DENSITY_FLOOR +#endif +#ifdef DENSITY_FLOOR } else if (strcmp(name, "density_floor") == 0) { parms->density_floor = atof(value); -# endif -# ifdef SCALAR_FLOOR +#endif +#ifdef SCALAR_FLOOR } else if (strcmp(name, "scalar_floor") == 0) { parms->scalar_floor = atof(value); -# endif +#endif #ifdef ANALYSIS } else if (strcmp(name, "analysis_scale_outputs_file") == 0) { strncpy(parms->analysis_scale_outputs_file, value, MAXLEN); diff --git a/src/gravity/gravity_functions.cpp b/src/gravity/gravity_functions.cpp index b92d06564..6c7a6dde7 100644 --- a/src/gravity/gravity_functions.cpp +++ b/src/gravity/gravity_functions.cpp @@ -137,7 +137,7 @@ void Grid3D::set_dt_Gravity() dt_particles = Calc_Particles_dt(); dt_particles = fmin(dt_particles, Particles.max_dt); #ifdef ONLY_PARTICLES - dt_min = dt_particles; + dt_min = dt_particles; chprintf(" dt_particles: %f \n", dt_particles); #else chprintf(" dt_hydro: %f dt_particles: %f \n", dt_hydro, dt_particles); @@ -211,7 +211,7 @@ Real Grav3D::Get_Average_Density() #ifdef MPI_CHOLLA dens_avrg_all = ReduceRealAvg(dens_mean); #else - dens_avrg_all = dens_mean; + dens_avrg_all = dens_mean; #endif dens_avrg = dens_avrg_all; @@ -530,8 +530,8 @@ void Grid3D::Compute_Gravitational_Potential(struct Parameters *P) input_density = Grav.F.density_d; output_potential = Grav.F.potential_d; #else - input_density = Grav.F.density_h; - output_potential = Grav.F.potential_h; + input_density = Grav.F.density_h; + output_potential = Grav.F.potential_h; #endif #ifdef SOR diff --git a/src/gravity/gravity_functions_gpu.cu b/src/gravity/gravity_functions_gpu.cu index b92d19084..15de64a95 100644 --- a/src/gravity/gravity_functions_gpu.cu +++ b/src/gravity/gravity_functions_gpu.cu @@ -127,7 +127,7 @@ void Grid3D::Copy_Hydro_Density_to_Gravity_GPU() #ifdef COSMOLOGY cosmo_rho_0_gas = Cosmo.rho_0_gas; #else - cosmo_rho_0_gas = 1.0; + cosmo_rho_0_gas = 1.0; #endif // Copy the density from the device array to the Poisson input density array @@ -261,7 +261,7 @@ void Grid3D::Extrapolate_Grav_Potential_GPU() #ifdef COSMOLOGY cosmo_factor = Cosmo.current_a * Cosmo.current_a / Cosmo.phi_0_gas; #else - cosmo_factor = 1.0; + cosmo_factor = 1.0; #endif // set values for GPU kernels diff --git a/src/grid/grid3D.cpp b/src/grid/grid3D.cpp index d844784d5..93499fea0 100644 --- a/src/grid/grid3D.cpp +++ b/src/grid/grid3D.cpp @@ -158,7 +158,7 @@ void Grid3D::Initialize(struct Parameters *P) #ifdef AVERAGE_SLOW_CELLS H.min_dt_slow = 1e-5; // Initialize the minumum dt to a tiny number -#endif // AVERAGE_SLOW_CELLS +#endif // AVERAGE_SLOW_CELLS #ifndef MPI_CHOLLA @@ -267,13 +267,13 @@ void Grid3D::Initialize(struct Parameters *P) #ifdef DENSITY_FLOOR H.density_floor = P->density_floor; #else - H.density_floor = 0.0; + H.density_floor = 0.0; #endif #ifdef SCALAR_FLOOR H.scalar_floor = P->scalar_floor; #else - H.scalar_floor = 0.0; + H.scalar_floor = 0.0; #endif #ifdef COSMOLOGY @@ -345,8 +345,8 @@ void Grid3D::AllocateMemory(void) GPU_Error_Check(cudaHostAlloc(&C.Grav_potential, H.n_cells * sizeof(Real), cudaHostAllocDefault)); GPU_Error_Check(cudaMalloc((void **)&C.d_Grav_potential, H.n_cells * sizeof(Real))); #else - C.Grav_potential = NULL; - C.d_Grav_potential = NULL; + C.Grav_potential = NULL; + C.d_Grav_potential = NULL; #endif #ifdef CHEMISTRY_GPU diff --git a/src/grid/initial_conditions.cpp b/src/grid/initial_conditions.cpp index 0e5aabc90..30e3eb459 100644 --- a/src/grid/initial_conditions.cpp +++ b/src/grid/initial_conditions.cpp @@ -1415,10 +1415,10 @@ void Grid3D::Clouds() C.GasEnergy[id] = p_cl / (gama - 1.0); #endif // DE #ifdef SCALAR -#ifdef DUST + #ifdef DUST C.host[id + H.n_cells * grid_enum::dust_density] = rho_cl * 1e-2; -#endif // DUST -#endif // SCALAR + #endif // DUST +#endif // SCALAR } } } @@ -1482,18 +1482,18 @@ void Grid3D::Zeldovich_Pancake(struct Parameters P) Real H0, h, Omega_M, rho_0, G, z_zeldovich, z_init, x_center, T_init, k_x; chprintf("Setting Zeldovich Pancake initial conditions...\n"); - H0 = P.H0; - h = H0 / 100; + H0 = P.H0; + h = H0 / 100; Omega_M = P.Omega_M; chprintf(" h = %f \n", h); chprintf(" Omega_M = %f \n", Omega_M); H0 /= 1000; //[km/s / kpc] - G = G_COSMO; - rho_0 = 3 * H0 * H0 / (8 * M_PI * G) * Omega_M / h / h; + G = G_COSMO; + rho_0 = 3 * H0 * H0 / (8 * M_PI * G) * Omega_M / h / h; z_zeldovich = 1; - z_init = P.Init_redshift; + z_init = P.Init_redshift; chprintf(" rho_0 = %f \n", rho_0); chprintf(" z_init = %f \n", z_init); chprintf(" z_zeldovich = %f \n", z_zeldovich); @@ -1553,17 +1553,17 @@ void Grid3D::Zeldovich_Pancake(struct Parameters P) index = (int(x_pos / H.dx) + 0) % 256; // index = ( index + 16 ) % 256; dens = ics_values[0 * nPoints + index]; - vel = ics_values[1 * nPoints + index]; - E = ics_values[2 * nPoints + index]; - U = ics_values[3 * nPoints + index]; + vel = ics_values[1 * nPoints + index]; + E = ics_values[2 * nPoints + index]; + U = ics_values[3 * nPoints + index]; // // // chprintf( "%f \n", vel ); - C.density[id] = dens; + C.density[id] = dens; C.momentum_x[id] = dens * vel; C.momentum_y[id] = 0; C.momentum_z[id] = 0; - C.Energy[id] = E; + C.Energy[id] = E; #ifdef DE C.GasEnergy[id] = U; diff --git a/src/integrators/VL_3D_cuda.cu b/src/integrators/VL_3D_cuda.cu index 25462fa91..2b4162ffc 100644 --- a/src/integrators/VL_3D_cuda.cu +++ b/src/integrators/VL_3D_cuda.cu @@ -32,7 +32,8 @@ __global__ void Update_Conserved_Variables_3D_half(Real *dev_conserved, Real *dev_conserved_half, Real *dev_F_x, Real *dev_F_y, Real *dev_F_z, int nx, int ny, int nz, int n_ghost, - Real dx, Real dy, Real dz, Real dt, Real gamma, int n_fields, Real density_floor); + Real dx, Real dy, Real dz, Real dt, Real gamma, int n_fields, + Real density_floor); 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, @@ -122,7 +123,7 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int #if defined(GRAVITY) dev_grav_potential = d_grav_potential; #else // not GRAVITY - dev_grav_potential = NULL; + dev_grav_potential = NULL; #endif // GRAVITY // If memory is single allocated: memory_allocated becomes true and @@ -327,11 +328,11 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int #endif // TEMPERATURE_FLOOR #ifdef SCALAR_FLOOR - #ifdef DUST + #ifdef DUST hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, grid_enum::dust_density, scalar_floor); GPU_Error_Check(); - #endif + #endif #endif // SCALAR_FLOOR return; diff --git a/src/integrators/simple_3D_cuda.cu b/src/integrators/simple_3D_cuda.cu index d3beb14b2..b68f0a351 100644 --- a/src/integrators/simple_3D_cuda.cu +++ b/src/integrators/simple_3D_cuda.cu @@ -187,11 +187,11 @@ void Simple_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, #endif // TEMPERATURE_FLOOR #ifdef SCALAR_FLOOR - #ifdef DUST + #ifdef DUST hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, grid_enum::dust_density, scalar_floor); CudaCheckError(); - #endif DUST + #endif DUST #endif // SCALAR_FLOOR return; diff --git a/src/io/io.cpp b/src/io/io.cpp index 9959267de..b243440e4 100644 --- a/src/io/io.cpp +++ b/src/io/io.cpp @@ -1400,12 +1400,12 @@ void Grid3D::Write_Grid_HDF5(hid_t file_id) #ifdef OUTPUT_METALS output_metals = true; #else // not OUTPUT_METALS - output_metals = false; + output_metals = false; #endif // OUTPUT_METALS #ifdef OUTPUT_ELECTRONS output_electrons = true; #else // not OUTPUT_ELECTRONS - output_electrons = false; + output_electrons = false; #endif // OUTPUT_ELECTRONS #ifdef OUTPUT_FULL_IONIZATION output_full_ionization = true; diff --git a/src/particles/io_particles.cpp b/src/particles/io_particles.cpp index e986c5287..e6da774ed 100644 --- a/src/particles/io_particles.cpp +++ b/src/particles/io_particles.cpp @@ -445,12 +445,12 @@ void Particles3D::Load_Particles_Data_HDF5(hid_t file_id, int nfile, struct Para Real vy_max_g = vy_max; Real vz_max_g = vz_max; - Real px_min_g = px_min; - Real py_min_g = py_min; - Real pz_min_g = pz_min; - Real vx_min_g = vx_min; - Real vy_min_g = vy_min; - Real vz_min_g = vz_min; + Real px_min_g = px_min; + Real py_min_g = py_min; + Real pz_min_g = pz_min; + Real vx_min_g = vx_min; + Real vy_min_g = vy_min; + Real vz_min_g = vz_min; #endif // MPI_CHOLLA // Print initial Statistics @@ -563,7 +563,7 @@ void Grid3D::Write_Particles_Data_HDF5(hid_t file_id) #ifdef MPI_CHOLLA N_particles_total = ReducePartIntSum(Particles.n_local); #else - N_particles_total = Particles.n_local; + N_particles_total = Particles.n_local; #endif // Print the total particles when saving the particles data diff --git a/src/particles/particles_3D.cpp b/src/particles/particles_3D.cpp index 6417e4136..87a2be8e5 100644 --- a/src/particles/particles_3D.cpp +++ b/src/particles/particles_3D.cpp @@ -157,12 +157,12 @@ void Particles3D::Initialize(struct Parameters *P, Grav3D &Grav, Real xbound, Re G.boundary_type_z0 = P->zlg_bcnd; G.boundary_type_z1 = P->zug_bcnd; #else - G.boundary_type_x0 = P->xl_bcnd; - G.boundary_type_x1 = P->xu_bcnd; - G.boundary_type_y0 = P->yl_bcnd; - G.boundary_type_y1 = P->yu_bcnd; - G.boundary_type_z0 = P->zl_bcnd; - G.boundary_type_z1 = P->zu_bcnd; + G.boundary_type_x0 = P->xl_bcnd; + G.boundary_type_x1 = P->xu_bcnd; + G.boundary_type_y0 = P->yl_bcnd; + G.boundary_type_y1 = P->yu_bcnd; + G.boundary_type_z0 = P->zl_bcnd; + G.boundary_type_z1 = P->zu_bcnd; #endif #ifdef PARTICLES_GPU @@ -211,7 +211,7 @@ void Particles3D::Initialize(struct Parameters *P, Grav3D &Grav, Real xbound, Re #ifdef MPI_CHOLLA n_total_initial = ReducePartIntSum(n_local); #else - n_total_initial = n_local; + n_total_initial = n_local; #endif chprintf("Particles Initialized: \n n_local: %lu \n", n_local); diff --git a/src/particles/particles_boundaries_cpu.cpp b/src/particles/particles_boundaries_cpu.cpp index 27470befe..772153534 100644 --- a/src/particles/particles_boundaries_cpu.cpp +++ b/src/particles/particles_boundaries_cpu.cpp @@ -433,13 +433,13 @@ void Particles3D::Unload_Particles_from_Buffer_CPU(int direction, int side, Real offset_extra += 1; pId = recv_buffer[offset_extra]; #else - pId = 0; + pId = 0; #endif #ifdef PARTICLE_AGE offset_extra += 1; pAge = recv_buffer[offset_extra]; #else - pAge = 0.0; + pAge = 0.0; #endif offset_buff += N_DATA_PER_PARTICLE_TRANSFER; diff --git a/src/reconstruction/plmp_cuda.cu b/src/reconstruction/plmp_cuda.cu index a000da4da..f69bbdc4b 100644 --- a/src/reconstruction/plmp_cuda.cu +++ b/src/reconstruction/plmp_cuda.cu @@ -120,7 +120,7 @@ __global__ void PLMP_cuda(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bou dge = dev_conserved[(n_fields - 1) * n_cells + id]; p_i = hydro_utilities::Get_Pressure_From_DE(E, E - E_kin, dge, gamma); #else - p_i = (dev_conserved[4 * n_cells + id] - 0.5 * d_i * (vx_i * vx_i + vy_i * vy_i + vz_i * vz_i)) * (gamma - 1.0); + p_i = (dev_conserved[4 * n_cells + id] - 0.5 * d_i * (vx_i * vx_i + vy_i * vy_i + vz_i * vz_i)) * (gamma - 1.0); #endif // PRESSURE_DE p_i = fmax(p_i, (Real)TINY_NUMBER); #ifdef SCALAR diff --git a/src/reconstruction/ppmc_cuda_tests.cu b/src/reconstruction/ppmc_cuda_tests.cu index 9e9b11140..c1319ea58 100644 --- a/src/reconstruction/ppmc_cuda_tests.cu +++ b/src/reconstruction/ppmc_cuda_tests.cu @@ -139,7 +139,7 @@ TEST(tALLPpmcVLReconstructor, CorrectInputExpectCorrectOutput) #ifdef MHD size_t const n_fields = 8; #else // not MHD - size_t const n_fields = 5; + size_t const n_fields = 5; #endif // MHD // Setup host grid. Fill host grid with random values and randomly assign maximum value diff --git a/src/reconstruction/ppmp_cuda.cu b/src/reconstruction/ppmp_cuda.cu index ae8da90cb..f84946437 100644 --- a/src/reconstruction/ppmp_cuda.cu +++ b/src/reconstruction/ppmp_cuda.cu @@ -166,7 +166,7 @@ __global__ void PPMP_cuda(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bou dge = dev_conserved[(n_fields - 1) * n_cells + id]; p_i = hydro_utilities::Get_Pressure_From_DE(E, E - E_kin, dge, gamma); #else - p_i = (dev_conserved[4 * n_cells + id] - 0.5 * d_i * (vx_i * vx_i + vy_i * vy_i + vz_i * vz_i)) * (gamma - 1.0); + p_i = (dev_conserved[4 * n_cells + id] - 0.5 * d_i * (vx_i * vx_i + vy_i * vy_i + vz_i * vz_i)) * (gamma - 1.0); #endif // PRESSURE_DE p_i = fmax(p_i, (Real)TINY_NUMBER); #ifdef DE diff --git a/src/reconstruction/reconstruction_tests.cu b/src/reconstruction/reconstruction_tests.cu index 6c2e19af7..74c0e6896 100644 --- a/src/reconstruction/reconstruction_tests.cu +++ b/src/reconstruction/reconstruction_tests.cu @@ -575,9 +575,13 @@ TEST(tALLReconstructionWriteData, CorrectInputExpectCorrectOutput) { // Set up test and mock up grid #ifdef MHD - reconstruction::Primitive interface{1, 2, 3, 4, 5, 6, 7, 8}; + reconstruction::Primitive interface { + 1, 2, 3, 4, 5, 6, 7, 8 + }; #else // MHD - reconstruction::Primitive interface{6, 7, 8, 9, 10}; + reconstruction::Primitive interface { + 6, 7, 8, 9, 10 + }; #endif // MHD size_t const nx = 3, ny = 3, nz = 3; size_t const n_cells = nx * ny * nz; diff --git a/src/system_tests/hydro_system_tests.cpp b/src/system_tests/hydro_system_tests.cpp index 6cffe9c21..18b2994bb 100644 --- a/src/system_tests/hydro_system_tests.cpp +++ b/src/system_tests/hydro_system_tests.cpp @@ -56,8 +56,8 @@ TEST_P(tHYDROtMHDSYSTEMSodShockTubeParameterizedMpi, CorrectInputExpectCorrectOu double const maxAllowedL1Error = 7.0E-3; double const maxAllowedError = 4.6E-2; #else - double const maxAllowedL1Error = 9.4E-5; - double const maxAllowedError = 6.4E-4; + double const maxAllowedL1Error = 9.4E-5; + double const maxAllowedError = 6.4E-4; #endif // MHD sodTest.numMpiRanks = GetParam(); From 399921b7e3fb4dd9835f008eb795f0027f319eb8 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Tue, 23 Jan 2024 11:50:19 -0500 Subject: [PATCH 20/40] remove reference to DENS_FLOOR macro --- src/particles/feedback_CIC_gpu.cu | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/src/particles/feedback_CIC_gpu.cu b/src/particles/feedback_CIC_gpu.cu index 0a4e8b292..bd162e585 100644 --- a/src/particles/feedback_CIC_gpu.cu +++ b/src/particles/feedback_CIC_gpu.cu @@ -157,9 +157,9 @@ __device__ Real GetSNRate(Real t, Real* dev_snr, Real snr_dt, Real t_start, Real } __device__ Real Calc_Timestep(Real gamma, Real* density, Real* momentum_x, Real* momentum_y, Real* momentum_z, - Real* energy, int index, Real dx, Real dy, Real dz) + Real* energy, int index, Real dx, Real dy, Real dz, Real density_floor) { - Real dens = fmax(density[index], DENS_FLOOR); + Real dens = fmax(density[index], density_floor); Real d_inv = 1.0 / dens; Real vx = momentum_x[index] * d_inv; Real vy = momentum_y[index] * d_inv; @@ -236,7 +236,7 @@ __global__ void Cluster_Feedback_Kernel(part_int_t n_local, part_int_t* id, Real Real* gasEnergy, Real* energy, Real* momentum_x, Real* momentum_y, Real* momentum_z, Real gamma, FeedbackPrng* states, Real* prev_dens, int* prev_N, short direction, Real* dev_snr, Real snr_dt, Real time_sn_start, - Real time_sn_end, int n_step) + Real time_sn_end, int n_step, Real density_floor) { __shared__ Real s_info[FEED_INFO_N * TPB_FEEDBACK]; // for collecting SN feedback information, like # // of SNe or # resolved. @@ -443,7 +443,7 @@ __global__ void Cluster_Feedback_Kernel(part_int_t n_local, part_int_t* id, Real if (direction > 0) { local_dti = fmax(local_dti, Calc_Timestep(gamma, density, momentum_x, momentum_y, momentum_z, - energy, indx, dx, dy, dz)); + energy, indx, dx, dy, dz, density_floor)); } } } @@ -605,7 +605,7 @@ __global__ void Cluster_Feedback_Kernel(part_int_t n_local, part_int_t* id, Real // kernel_printf("urs time:%.3e id:%d N:%d d:%.5e\n", t, // id[gtid], N, n_0); local_dti = fmax(local_dti, Calc_Timestep(gamma, density, momentum_x, momentum_y, momentum_z, - energy, indx, dx, dy, dz)); + energy, indx, dx, dy, dz, density_floor)); } } } @@ -698,7 +698,7 @@ Real supernova::Cluster_Feedback(Grid3D& G, FeedbackAnalysis& analysis) G.H.nz, G.H.n_ghost, G.H.t, G.H.dt, d_dti, d_info, G.C.d_density, G.C.d_GasEnergy, G.C.d_Energy, G.C.d_momentum_x, G.C.d_momentum_y, G.C.d_momentum_z, gama, supernova::randStates, d_prev_dens, d_prev_N, direction, dev_snr, snr_dt, time_sn_start, - time_sn_end, G.H.n_step); + time_sn_end, G.H.n_step, G.H.density_floor); GPU_Error_Check(cudaMemcpy(&h_dti, d_dti, sizeof(Real), cudaMemcpyDeviceToHost)); } @@ -719,7 +719,7 @@ Real supernova::Cluster_Feedback(Grid3D& G, FeedbackAnalysis& analysis) G.H.nz, G.H.n_ghost, G.H.t, G.H.dt, d_dti, d_info, G.C.d_density, G.C.d_GasEnergy, G.C.d_Energy, G.C.d_momentum_x, G.C.d_momentum_y, G.C.d_momentum_z, gama, supernova::randStates, d_prev_dens, d_prev_N, direction, dev_snr, snr_dt, time_sn_start, - time_sn_end, G.H.n_step); + time_sn_end, G.H.n_step, G.H.density_floor); GPU_Error_Check(cudaDeviceSynchronize()); } From f276258598f8997163907907b032672ad87e346e Mon Sep 17 00:00:00 2001 From: helenarichie Date: Tue, 23 Jan 2024 16:21:26 -0500 Subject: [PATCH 21/40] run clang format --- src/analysis/feedback_analysis.cpp | 2 +- src/chemistry_gpu/chemistry_functions.cpp | 2 +- src/cooling_grackle/cool_grackle.cpp | 2 +- src/gravity/gravity_functions.cpp | 8 ++++---- src/gravity/gravity_functions_gpu.cu | 4 ++-- src/grid/grid3D.cpp | 8 ++++---- src/grid/initial_conditions.cpp | 20 ++++++++++---------- src/integrators/VL_3D_cuda.cu | 2 +- src/io/io.cpp | 4 ++-- src/particles/io_particles.cpp | 14 +++++++------- src/particles/particles_3D.cpp | 14 +++++++------- src/particles/particles_boundaries_cpu.cpp | 4 ++-- src/reconstruction/plmp_cuda.cu | 2 +- src/reconstruction/ppmc_cuda_tests.cu | 2 +- src/reconstruction/ppmp_cuda.cu | 2 +- src/reconstruction/reconstruction_tests.cu | 8 ++------ src/system_tests/hydro_system_tests.cpp | 4 ++-- 17 files changed, 49 insertions(+), 53 deletions(-) diff --git a/src/analysis/feedback_analysis.cpp b/src/analysis/feedback_analysis.cpp index 5fe0e7543..3dab7b6da 100644 --- a/src/analysis/feedback_analysis.cpp +++ b/src/analysis/feedback_analysis.cpp @@ -87,7 +87,7 @@ void FeedbackAnalysis::Compute_Gas_Velocity_Dispersion(Grid3D& G) #ifdef MPI_CHOLLA MPI_Allreduce(&partial_mass, &total_mass, 1, MPI_CHREAL, MPI_SUM, world); #else - total_mass = partial_mass; + total_mass = partial_mass; #endif for (k = G.H.n_ghost; k < G.H.nz - G.H.n_ghost; k++) { diff --git a/src/chemistry_gpu/chemistry_functions.cpp b/src/chemistry_gpu/chemistry_functions.cpp index 65e3af691..7999a6d55 100644 --- a/src/chemistry_gpu/chemistry_functions.cpp +++ b/src/chemistry_gpu/chemistry_functions.cpp @@ -228,7 +228,7 @@ void Grid3D::Update_Chemistry() #ifdef COSMOLOGY Chem.H.current_z = Cosmo.current_z; #else - Chem.H.current_z = 0; + Chem.H.current_z = 0; #endif Do_Chemistry_Update(C.device, H.nx, H.ny, H.nz, H.n_ghost, H.n_fields, H.dt, Chem.H); diff --git a/src/cooling_grackle/cool_grackle.cpp b/src/cooling_grackle/cool_grackle.cpp index c5f2a8078..a7f5c36cb 100644 --- a/src/cooling_grackle/cool_grackle.cpp +++ b/src/cooling_grackle/cool_grackle.cpp @@ -89,7 +89,7 @@ void Cool_GK::Initialize(struct Parameters *P, Cosmology &Cosmo) data->metal_cooling = 1; // metal cooling off #else chprintf("WARNING: Metal Cooling is Off. \n"); - data->metal_cooling = 0; // metal cooling off + data->metal_cooling = 0; // metal cooling off #endif #ifdef PARALLEL_OMP diff --git a/src/gravity/gravity_functions.cpp b/src/gravity/gravity_functions.cpp index 6c7a6dde7..b92d06564 100644 --- a/src/gravity/gravity_functions.cpp +++ b/src/gravity/gravity_functions.cpp @@ -137,7 +137,7 @@ void Grid3D::set_dt_Gravity() dt_particles = Calc_Particles_dt(); dt_particles = fmin(dt_particles, Particles.max_dt); #ifdef ONLY_PARTICLES - dt_min = dt_particles; + dt_min = dt_particles; chprintf(" dt_particles: %f \n", dt_particles); #else chprintf(" dt_hydro: %f dt_particles: %f \n", dt_hydro, dt_particles); @@ -211,7 +211,7 @@ Real Grav3D::Get_Average_Density() #ifdef MPI_CHOLLA dens_avrg_all = ReduceRealAvg(dens_mean); #else - dens_avrg_all = dens_mean; + dens_avrg_all = dens_mean; #endif dens_avrg = dens_avrg_all; @@ -530,8 +530,8 @@ void Grid3D::Compute_Gravitational_Potential(struct Parameters *P) input_density = Grav.F.density_d; output_potential = Grav.F.potential_d; #else - input_density = Grav.F.density_h; - output_potential = Grav.F.potential_h; + input_density = Grav.F.density_h; + output_potential = Grav.F.potential_h; #endif #ifdef SOR diff --git a/src/gravity/gravity_functions_gpu.cu b/src/gravity/gravity_functions_gpu.cu index 15de64a95..b92d19084 100644 --- a/src/gravity/gravity_functions_gpu.cu +++ b/src/gravity/gravity_functions_gpu.cu @@ -127,7 +127,7 @@ void Grid3D::Copy_Hydro_Density_to_Gravity_GPU() #ifdef COSMOLOGY cosmo_rho_0_gas = Cosmo.rho_0_gas; #else - cosmo_rho_0_gas = 1.0; + cosmo_rho_0_gas = 1.0; #endif // Copy the density from the device array to the Poisson input density array @@ -261,7 +261,7 @@ void Grid3D::Extrapolate_Grav_Potential_GPU() #ifdef COSMOLOGY cosmo_factor = Cosmo.current_a * Cosmo.current_a / Cosmo.phi_0_gas; #else - cosmo_factor = 1.0; + cosmo_factor = 1.0; #endif // set values for GPU kernels diff --git a/src/grid/grid3D.cpp b/src/grid/grid3D.cpp index 93499fea0..dd65da756 100644 --- a/src/grid/grid3D.cpp +++ b/src/grid/grid3D.cpp @@ -267,13 +267,13 @@ void Grid3D::Initialize(struct Parameters *P) #ifdef DENSITY_FLOOR H.density_floor = P->density_floor; #else - H.density_floor = 0.0; + H.density_floor = 0.0; #endif #ifdef SCALAR_FLOOR H.scalar_floor = P->scalar_floor; #else - H.scalar_floor = 0.0; + H.scalar_floor = 0.0; #endif #ifdef COSMOLOGY @@ -345,8 +345,8 @@ void Grid3D::AllocateMemory(void) GPU_Error_Check(cudaHostAlloc(&C.Grav_potential, H.n_cells * sizeof(Real), cudaHostAllocDefault)); GPU_Error_Check(cudaMalloc((void **)&C.d_Grav_potential, H.n_cells * sizeof(Real))); #else - C.Grav_potential = NULL; - C.d_Grav_potential = NULL; + C.Grav_potential = NULL; + C.d_Grav_potential = NULL; #endif #ifdef CHEMISTRY_GPU diff --git a/src/grid/initial_conditions.cpp b/src/grid/initial_conditions.cpp index 30e3eb459..af558be8f 100644 --- a/src/grid/initial_conditions.cpp +++ b/src/grid/initial_conditions.cpp @@ -1482,18 +1482,18 @@ void Grid3D::Zeldovich_Pancake(struct Parameters P) Real H0, h, Omega_M, rho_0, G, z_zeldovich, z_init, x_center, T_init, k_x; chprintf("Setting Zeldovich Pancake initial conditions...\n"); - H0 = P.H0; - h = H0 / 100; + H0 = P.H0; + h = H0 / 100; Omega_M = P.Omega_M; chprintf(" h = %f \n", h); chprintf(" Omega_M = %f \n", Omega_M); H0 /= 1000; //[km/s / kpc] - G = G_COSMO; - rho_0 = 3 * H0 * H0 / (8 * M_PI * G) * Omega_M / h / h; + G = G_COSMO; + rho_0 = 3 * H0 * H0 / (8 * M_PI * G) * Omega_M / h / h; z_zeldovich = 1; - z_init = P.Init_redshift; + z_init = P.Init_redshift; chprintf(" rho_0 = %f \n", rho_0); chprintf(" z_init = %f \n", z_init); chprintf(" z_zeldovich = %f \n", z_zeldovich); @@ -1553,17 +1553,17 @@ void Grid3D::Zeldovich_Pancake(struct Parameters P) index = (int(x_pos / H.dx) + 0) % 256; // index = ( index + 16 ) % 256; dens = ics_values[0 * nPoints + index]; - vel = ics_values[1 * nPoints + index]; - E = ics_values[2 * nPoints + index]; - U = ics_values[3 * nPoints + index]; + vel = ics_values[1 * nPoints + index]; + E = ics_values[2 * nPoints + index]; + U = ics_values[3 * nPoints + index]; // // // chprintf( "%f \n", vel ); - C.density[id] = dens; + C.density[id] = dens; C.momentum_x[id] = dens * vel; C.momentum_y[id] = 0; C.momentum_z[id] = 0; - C.Energy[id] = E; + C.Energy[id] = E; #ifdef DE C.GasEnergy[id] = U; diff --git a/src/integrators/VL_3D_cuda.cu b/src/integrators/VL_3D_cuda.cu index 2b4162ffc..1118c1419 100644 --- a/src/integrators/VL_3D_cuda.cu +++ b/src/integrators/VL_3D_cuda.cu @@ -123,7 +123,7 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int #if defined(GRAVITY) dev_grav_potential = d_grav_potential; #else // not GRAVITY - dev_grav_potential = NULL; + dev_grav_potential = NULL; #endif // GRAVITY // If memory is single allocated: memory_allocated becomes true and diff --git a/src/io/io.cpp b/src/io/io.cpp index b243440e4..9959267de 100644 --- a/src/io/io.cpp +++ b/src/io/io.cpp @@ -1400,12 +1400,12 @@ void Grid3D::Write_Grid_HDF5(hid_t file_id) #ifdef OUTPUT_METALS output_metals = true; #else // not OUTPUT_METALS - output_metals = false; + output_metals = false; #endif // OUTPUT_METALS #ifdef OUTPUT_ELECTRONS output_electrons = true; #else // not OUTPUT_ELECTRONS - output_electrons = false; + output_electrons = false; #endif // OUTPUT_ELECTRONS #ifdef OUTPUT_FULL_IONIZATION output_full_ionization = true; diff --git a/src/particles/io_particles.cpp b/src/particles/io_particles.cpp index e6da774ed..e986c5287 100644 --- a/src/particles/io_particles.cpp +++ b/src/particles/io_particles.cpp @@ -445,12 +445,12 @@ void Particles3D::Load_Particles_Data_HDF5(hid_t file_id, int nfile, struct Para Real vy_max_g = vy_max; Real vz_max_g = vz_max; - Real px_min_g = px_min; - Real py_min_g = py_min; - Real pz_min_g = pz_min; - Real vx_min_g = vx_min; - Real vy_min_g = vy_min; - Real vz_min_g = vz_min; + Real px_min_g = px_min; + Real py_min_g = py_min; + Real pz_min_g = pz_min; + Real vx_min_g = vx_min; + Real vy_min_g = vy_min; + Real vz_min_g = vz_min; #endif // MPI_CHOLLA // Print initial Statistics @@ -563,7 +563,7 @@ void Grid3D::Write_Particles_Data_HDF5(hid_t file_id) #ifdef MPI_CHOLLA N_particles_total = ReducePartIntSum(Particles.n_local); #else - N_particles_total = Particles.n_local; + N_particles_total = Particles.n_local; #endif // Print the total particles when saving the particles data diff --git a/src/particles/particles_3D.cpp b/src/particles/particles_3D.cpp index 87a2be8e5..6417e4136 100644 --- a/src/particles/particles_3D.cpp +++ b/src/particles/particles_3D.cpp @@ -157,12 +157,12 @@ void Particles3D::Initialize(struct Parameters *P, Grav3D &Grav, Real xbound, Re G.boundary_type_z0 = P->zlg_bcnd; G.boundary_type_z1 = P->zug_bcnd; #else - G.boundary_type_x0 = P->xl_bcnd; - G.boundary_type_x1 = P->xu_bcnd; - G.boundary_type_y0 = P->yl_bcnd; - G.boundary_type_y1 = P->yu_bcnd; - G.boundary_type_z0 = P->zl_bcnd; - G.boundary_type_z1 = P->zu_bcnd; + G.boundary_type_x0 = P->xl_bcnd; + G.boundary_type_x1 = P->xu_bcnd; + G.boundary_type_y0 = P->yl_bcnd; + G.boundary_type_y1 = P->yu_bcnd; + G.boundary_type_z0 = P->zl_bcnd; + G.boundary_type_z1 = P->zu_bcnd; #endif #ifdef PARTICLES_GPU @@ -211,7 +211,7 @@ void Particles3D::Initialize(struct Parameters *P, Grav3D &Grav, Real xbound, Re #ifdef MPI_CHOLLA n_total_initial = ReducePartIntSum(n_local); #else - n_total_initial = n_local; + n_total_initial = n_local; #endif chprintf("Particles Initialized: \n n_local: %lu \n", n_local); diff --git a/src/particles/particles_boundaries_cpu.cpp b/src/particles/particles_boundaries_cpu.cpp index 772153534..27470befe 100644 --- a/src/particles/particles_boundaries_cpu.cpp +++ b/src/particles/particles_boundaries_cpu.cpp @@ -433,13 +433,13 @@ void Particles3D::Unload_Particles_from_Buffer_CPU(int direction, int side, Real offset_extra += 1; pId = recv_buffer[offset_extra]; #else - pId = 0; + pId = 0; #endif #ifdef PARTICLE_AGE offset_extra += 1; pAge = recv_buffer[offset_extra]; #else - pAge = 0.0; + pAge = 0.0; #endif offset_buff += N_DATA_PER_PARTICLE_TRANSFER; diff --git a/src/reconstruction/plmp_cuda.cu b/src/reconstruction/plmp_cuda.cu index f69bbdc4b..a000da4da 100644 --- a/src/reconstruction/plmp_cuda.cu +++ b/src/reconstruction/plmp_cuda.cu @@ -120,7 +120,7 @@ __global__ void PLMP_cuda(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bou dge = dev_conserved[(n_fields - 1) * n_cells + id]; p_i = hydro_utilities::Get_Pressure_From_DE(E, E - E_kin, dge, gamma); #else - p_i = (dev_conserved[4 * n_cells + id] - 0.5 * d_i * (vx_i * vx_i + vy_i * vy_i + vz_i * vz_i)) * (gamma - 1.0); + p_i = (dev_conserved[4 * n_cells + id] - 0.5 * d_i * (vx_i * vx_i + vy_i * vy_i + vz_i * vz_i)) * (gamma - 1.0); #endif // PRESSURE_DE p_i = fmax(p_i, (Real)TINY_NUMBER); #ifdef SCALAR diff --git a/src/reconstruction/ppmc_cuda_tests.cu b/src/reconstruction/ppmc_cuda_tests.cu index c1319ea58..9e9b11140 100644 --- a/src/reconstruction/ppmc_cuda_tests.cu +++ b/src/reconstruction/ppmc_cuda_tests.cu @@ -139,7 +139,7 @@ TEST(tALLPpmcVLReconstructor, CorrectInputExpectCorrectOutput) #ifdef MHD size_t const n_fields = 8; #else // not MHD - size_t const n_fields = 5; + size_t const n_fields = 5; #endif // MHD // Setup host grid. Fill host grid with random values and randomly assign maximum value diff --git a/src/reconstruction/ppmp_cuda.cu b/src/reconstruction/ppmp_cuda.cu index f84946437..ae8da90cb 100644 --- a/src/reconstruction/ppmp_cuda.cu +++ b/src/reconstruction/ppmp_cuda.cu @@ -166,7 +166,7 @@ __global__ void PPMP_cuda(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bou dge = dev_conserved[(n_fields - 1) * n_cells + id]; p_i = hydro_utilities::Get_Pressure_From_DE(E, E - E_kin, dge, gamma); #else - p_i = (dev_conserved[4 * n_cells + id] - 0.5 * d_i * (vx_i * vx_i + vy_i * vy_i + vz_i * vz_i)) * (gamma - 1.0); + p_i = (dev_conserved[4 * n_cells + id] - 0.5 * d_i * (vx_i * vx_i + vy_i * vy_i + vz_i * vz_i)) * (gamma - 1.0); #endif // PRESSURE_DE p_i = fmax(p_i, (Real)TINY_NUMBER); #ifdef DE diff --git a/src/reconstruction/reconstruction_tests.cu b/src/reconstruction/reconstruction_tests.cu index 74c0e6896..6c2e19af7 100644 --- a/src/reconstruction/reconstruction_tests.cu +++ b/src/reconstruction/reconstruction_tests.cu @@ -575,13 +575,9 @@ TEST(tALLReconstructionWriteData, CorrectInputExpectCorrectOutput) { // Set up test and mock up grid #ifdef MHD - reconstruction::Primitive interface { - 1, 2, 3, 4, 5, 6, 7, 8 - }; + reconstruction::Primitive interface{1, 2, 3, 4, 5, 6, 7, 8}; #else // MHD - reconstruction::Primitive interface { - 6, 7, 8, 9, 10 - }; + reconstruction::Primitive interface{6, 7, 8, 9, 10}; #endif // MHD size_t const nx = 3, ny = 3, nz = 3; size_t const n_cells = nx * ny * nz; diff --git a/src/system_tests/hydro_system_tests.cpp b/src/system_tests/hydro_system_tests.cpp index 18b2994bb..6cffe9c21 100644 --- a/src/system_tests/hydro_system_tests.cpp +++ b/src/system_tests/hydro_system_tests.cpp @@ -56,8 +56,8 @@ TEST_P(tHYDROtMHDSYSTEMSodShockTubeParameterizedMpi, CorrectInputExpectCorrectOu double const maxAllowedL1Error = 7.0E-3; double const maxAllowedError = 4.6E-2; #else - double const maxAllowedL1Error = 9.4E-5; - double const maxAllowedError = 6.4E-4; + double const maxAllowedL1Error = 9.4E-5; + double const maxAllowedError = 6.4E-4; #endif // MHD sodTest.numMpiRanks = GetParam(); From 37ccafdced47c8eccc07eb92822e2a0e97d119e8 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Tue, 23 Jan 2024 16:22:33 -0500 Subject: [PATCH 22/40] run clang format --- src/analysis/feedback_analysis.cpp | 2 +- src/chemistry_gpu/chemistry_functions.cpp | 2 +- src/cooling_grackle/cool_grackle.cpp | 2 +- src/gravity/gravity_functions.cpp | 8 ++++---- src/gravity/gravity_functions_gpu.cu | 4 ++-- src/grid/grid3D.cpp | 8 ++++---- src/grid/initial_conditions.cpp | 20 ++++++++++---------- src/integrators/VL_3D_cuda.cu | 2 +- src/io/io.cpp | 4 ++-- src/particles/io_particles.cpp | 14 +++++++------- src/particles/particles_3D.cpp | 14 +++++++------- src/particles/particles_boundaries_cpu.cpp | 4 ++-- src/reconstruction/plmp_cuda.cu | 2 +- src/reconstruction/ppmc_cuda_tests.cu | 2 +- src/reconstruction/ppmp_cuda.cu | 2 +- src/reconstruction/reconstruction_tests.cu | 8 ++++++-- src/system_tests/hydro_system_tests.cpp | 4 ++-- 17 files changed, 53 insertions(+), 49 deletions(-) diff --git a/src/analysis/feedback_analysis.cpp b/src/analysis/feedback_analysis.cpp index 3dab7b6da..5fe0e7543 100644 --- a/src/analysis/feedback_analysis.cpp +++ b/src/analysis/feedback_analysis.cpp @@ -87,7 +87,7 @@ void FeedbackAnalysis::Compute_Gas_Velocity_Dispersion(Grid3D& G) #ifdef MPI_CHOLLA MPI_Allreduce(&partial_mass, &total_mass, 1, MPI_CHREAL, MPI_SUM, world); #else - total_mass = partial_mass; + total_mass = partial_mass; #endif for (k = G.H.n_ghost; k < G.H.nz - G.H.n_ghost; k++) { diff --git a/src/chemistry_gpu/chemistry_functions.cpp b/src/chemistry_gpu/chemistry_functions.cpp index 7999a6d55..65e3af691 100644 --- a/src/chemistry_gpu/chemistry_functions.cpp +++ b/src/chemistry_gpu/chemistry_functions.cpp @@ -228,7 +228,7 @@ void Grid3D::Update_Chemistry() #ifdef COSMOLOGY Chem.H.current_z = Cosmo.current_z; #else - Chem.H.current_z = 0; + Chem.H.current_z = 0; #endif Do_Chemistry_Update(C.device, H.nx, H.ny, H.nz, H.n_ghost, H.n_fields, H.dt, Chem.H); diff --git a/src/cooling_grackle/cool_grackle.cpp b/src/cooling_grackle/cool_grackle.cpp index a7f5c36cb..c5f2a8078 100644 --- a/src/cooling_grackle/cool_grackle.cpp +++ b/src/cooling_grackle/cool_grackle.cpp @@ -89,7 +89,7 @@ void Cool_GK::Initialize(struct Parameters *P, Cosmology &Cosmo) data->metal_cooling = 1; // metal cooling off #else chprintf("WARNING: Metal Cooling is Off. \n"); - data->metal_cooling = 0; // metal cooling off + data->metal_cooling = 0; // metal cooling off #endif #ifdef PARALLEL_OMP diff --git a/src/gravity/gravity_functions.cpp b/src/gravity/gravity_functions.cpp index b92d06564..6c7a6dde7 100644 --- a/src/gravity/gravity_functions.cpp +++ b/src/gravity/gravity_functions.cpp @@ -137,7 +137,7 @@ void Grid3D::set_dt_Gravity() dt_particles = Calc_Particles_dt(); dt_particles = fmin(dt_particles, Particles.max_dt); #ifdef ONLY_PARTICLES - dt_min = dt_particles; + dt_min = dt_particles; chprintf(" dt_particles: %f \n", dt_particles); #else chprintf(" dt_hydro: %f dt_particles: %f \n", dt_hydro, dt_particles); @@ -211,7 +211,7 @@ Real Grav3D::Get_Average_Density() #ifdef MPI_CHOLLA dens_avrg_all = ReduceRealAvg(dens_mean); #else - dens_avrg_all = dens_mean; + dens_avrg_all = dens_mean; #endif dens_avrg = dens_avrg_all; @@ -530,8 +530,8 @@ void Grid3D::Compute_Gravitational_Potential(struct Parameters *P) input_density = Grav.F.density_d; output_potential = Grav.F.potential_d; #else - input_density = Grav.F.density_h; - output_potential = Grav.F.potential_h; + input_density = Grav.F.density_h; + output_potential = Grav.F.potential_h; #endif #ifdef SOR diff --git a/src/gravity/gravity_functions_gpu.cu b/src/gravity/gravity_functions_gpu.cu index b92d19084..15de64a95 100644 --- a/src/gravity/gravity_functions_gpu.cu +++ b/src/gravity/gravity_functions_gpu.cu @@ -127,7 +127,7 @@ void Grid3D::Copy_Hydro_Density_to_Gravity_GPU() #ifdef COSMOLOGY cosmo_rho_0_gas = Cosmo.rho_0_gas; #else - cosmo_rho_0_gas = 1.0; + cosmo_rho_0_gas = 1.0; #endif // Copy the density from the device array to the Poisson input density array @@ -261,7 +261,7 @@ void Grid3D::Extrapolate_Grav_Potential_GPU() #ifdef COSMOLOGY cosmo_factor = Cosmo.current_a * Cosmo.current_a / Cosmo.phi_0_gas; #else - cosmo_factor = 1.0; + cosmo_factor = 1.0; #endif // set values for GPU kernels diff --git a/src/grid/grid3D.cpp b/src/grid/grid3D.cpp index dd65da756..93499fea0 100644 --- a/src/grid/grid3D.cpp +++ b/src/grid/grid3D.cpp @@ -267,13 +267,13 @@ void Grid3D::Initialize(struct Parameters *P) #ifdef DENSITY_FLOOR H.density_floor = P->density_floor; #else - H.density_floor = 0.0; + H.density_floor = 0.0; #endif #ifdef SCALAR_FLOOR H.scalar_floor = P->scalar_floor; #else - H.scalar_floor = 0.0; + H.scalar_floor = 0.0; #endif #ifdef COSMOLOGY @@ -345,8 +345,8 @@ void Grid3D::AllocateMemory(void) GPU_Error_Check(cudaHostAlloc(&C.Grav_potential, H.n_cells * sizeof(Real), cudaHostAllocDefault)); GPU_Error_Check(cudaMalloc((void **)&C.d_Grav_potential, H.n_cells * sizeof(Real))); #else - C.Grav_potential = NULL; - C.d_Grav_potential = NULL; + C.Grav_potential = NULL; + C.d_Grav_potential = NULL; #endif #ifdef CHEMISTRY_GPU diff --git a/src/grid/initial_conditions.cpp b/src/grid/initial_conditions.cpp index af558be8f..30e3eb459 100644 --- a/src/grid/initial_conditions.cpp +++ b/src/grid/initial_conditions.cpp @@ -1482,18 +1482,18 @@ void Grid3D::Zeldovich_Pancake(struct Parameters P) Real H0, h, Omega_M, rho_0, G, z_zeldovich, z_init, x_center, T_init, k_x; chprintf("Setting Zeldovich Pancake initial conditions...\n"); - H0 = P.H0; - h = H0 / 100; + H0 = P.H0; + h = H0 / 100; Omega_M = P.Omega_M; chprintf(" h = %f \n", h); chprintf(" Omega_M = %f \n", Omega_M); H0 /= 1000; //[km/s / kpc] - G = G_COSMO; - rho_0 = 3 * H0 * H0 / (8 * M_PI * G) * Omega_M / h / h; + G = G_COSMO; + rho_0 = 3 * H0 * H0 / (8 * M_PI * G) * Omega_M / h / h; z_zeldovich = 1; - z_init = P.Init_redshift; + z_init = P.Init_redshift; chprintf(" rho_0 = %f \n", rho_0); chprintf(" z_init = %f \n", z_init); chprintf(" z_zeldovich = %f \n", z_zeldovich); @@ -1553,17 +1553,17 @@ void Grid3D::Zeldovich_Pancake(struct Parameters P) index = (int(x_pos / H.dx) + 0) % 256; // index = ( index + 16 ) % 256; dens = ics_values[0 * nPoints + index]; - vel = ics_values[1 * nPoints + index]; - E = ics_values[2 * nPoints + index]; - U = ics_values[3 * nPoints + index]; + vel = ics_values[1 * nPoints + index]; + E = ics_values[2 * nPoints + index]; + U = ics_values[3 * nPoints + index]; // // // chprintf( "%f \n", vel ); - C.density[id] = dens; + C.density[id] = dens; C.momentum_x[id] = dens * vel; C.momentum_y[id] = 0; C.momentum_z[id] = 0; - C.Energy[id] = E; + C.Energy[id] = E; #ifdef DE C.GasEnergy[id] = U; diff --git a/src/integrators/VL_3D_cuda.cu b/src/integrators/VL_3D_cuda.cu index 1118c1419..2b4162ffc 100644 --- a/src/integrators/VL_3D_cuda.cu +++ b/src/integrators/VL_3D_cuda.cu @@ -123,7 +123,7 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int #if defined(GRAVITY) dev_grav_potential = d_grav_potential; #else // not GRAVITY - dev_grav_potential = NULL; + dev_grav_potential = NULL; #endif // GRAVITY // If memory is single allocated: memory_allocated becomes true and diff --git a/src/io/io.cpp b/src/io/io.cpp index 9959267de..b243440e4 100644 --- a/src/io/io.cpp +++ b/src/io/io.cpp @@ -1400,12 +1400,12 @@ void Grid3D::Write_Grid_HDF5(hid_t file_id) #ifdef OUTPUT_METALS output_metals = true; #else // not OUTPUT_METALS - output_metals = false; + output_metals = false; #endif // OUTPUT_METALS #ifdef OUTPUT_ELECTRONS output_electrons = true; #else // not OUTPUT_ELECTRONS - output_electrons = false; + output_electrons = false; #endif // OUTPUT_ELECTRONS #ifdef OUTPUT_FULL_IONIZATION output_full_ionization = true; diff --git a/src/particles/io_particles.cpp b/src/particles/io_particles.cpp index e986c5287..e6da774ed 100644 --- a/src/particles/io_particles.cpp +++ b/src/particles/io_particles.cpp @@ -445,12 +445,12 @@ void Particles3D::Load_Particles_Data_HDF5(hid_t file_id, int nfile, struct Para Real vy_max_g = vy_max; Real vz_max_g = vz_max; - Real px_min_g = px_min; - Real py_min_g = py_min; - Real pz_min_g = pz_min; - Real vx_min_g = vx_min; - Real vy_min_g = vy_min; - Real vz_min_g = vz_min; + Real px_min_g = px_min; + Real py_min_g = py_min; + Real pz_min_g = pz_min; + Real vx_min_g = vx_min; + Real vy_min_g = vy_min; + Real vz_min_g = vz_min; #endif // MPI_CHOLLA // Print initial Statistics @@ -563,7 +563,7 @@ void Grid3D::Write_Particles_Data_HDF5(hid_t file_id) #ifdef MPI_CHOLLA N_particles_total = ReducePartIntSum(Particles.n_local); #else - N_particles_total = Particles.n_local; + N_particles_total = Particles.n_local; #endif // Print the total particles when saving the particles data diff --git a/src/particles/particles_3D.cpp b/src/particles/particles_3D.cpp index 6417e4136..87a2be8e5 100644 --- a/src/particles/particles_3D.cpp +++ b/src/particles/particles_3D.cpp @@ -157,12 +157,12 @@ void Particles3D::Initialize(struct Parameters *P, Grav3D &Grav, Real xbound, Re G.boundary_type_z0 = P->zlg_bcnd; G.boundary_type_z1 = P->zug_bcnd; #else - G.boundary_type_x0 = P->xl_bcnd; - G.boundary_type_x1 = P->xu_bcnd; - G.boundary_type_y0 = P->yl_bcnd; - G.boundary_type_y1 = P->yu_bcnd; - G.boundary_type_z0 = P->zl_bcnd; - G.boundary_type_z1 = P->zu_bcnd; + G.boundary_type_x0 = P->xl_bcnd; + G.boundary_type_x1 = P->xu_bcnd; + G.boundary_type_y0 = P->yl_bcnd; + G.boundary_type_y1 = P->yu_bcnd; + G.boundary_type_z0 = P->zl_bcnd; + G.boundary_type_z1 = P->zu_bcnd; #endif #ifdef PARTICLES_GPU @@ -211,7 +211,7 @@ void Particles3D::Initialize(struct Parameters *P, Grav3D &Grav, Real xbound, Re #ifdef MPI_CHOLLA n_total_initial = ReducePartIntSum(n_local); #else - n_total_initial = n_local; + n_total_initial = n_local; #endif chprintf("Particles Initialized: \n n_local: %lu \n", n_local); diff --git a/src/particles/particles_boundaries_cpu.cpp b/src/particles/particles_boundaries_cpu.cpp index 27470befe..772153534 100644 --- a/src/particles/particles_boundaries_cpu.cpp +++ b/src/particles/particles_boundaries_cpu.cpp @@ -433,13 +433,13 @@ void Particles3D::Unload_Particles_from_Buffer_CPU(int direction, int side, Real offset_extra += 1; pId = recv_buffer[offset_extra]; #else - pId = 0; + pId = 0; #endif #ifdef PARTICLE_AGE offset_extra += 1; pAge = recv_buffer[offset_extra]; #else - pAge = 0.0; + pAge = 0.0; #endif offset_buff += N_DATA_PER_PARTICLE_TRANSFER; diff --git a/src/reconstruction/plmp_cuda.cu b/src/reconstruction/plmp_cuda.cu index a000da4da..f69bbdc4b 100644 --- a/src/reconstruction/plmp_cuda.cu +++ b/src/reconstruction/plmp_cuda.cu @@ -120,7 +120,7 @@ __global__ void PLMP_cuda(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bou dge = dev_conserved[(n_fields - 1) * n_cells + id]; p_i = hydro_utilities::Get_Pressure_From_DE(E, E - E_kin, dge, gamma); #else - p_i = (dev_conserved[4 * n_cells + id] - 0.5 * d_i * (vx_i * vx_i + vy_i * vy_i + vz_i * vz_i)) * (gamma - 1.0); + p_i = (dev_conserved[4 * n_cells + id] - 0.5 * d_i * (vx_i * vx_i + vy_i * vy_i + vz_i * vz_i)) * (gamma - 1.0); #endif // PRESSURE_DE p_i = fmax(p_i, (Real)TINY_NUMBER); #ifdef SCALAR diff --git a/src/reconstruction/ppmc_cuda_tests.cu b/src/reconstruction/ppmc_cuda_tests.cu index 9e9b11140..c1319ea58 100644 --- a/src/reconstruction/ppmc_cuda_tests.cu +++ b/src/reconstruction/ppmc_cuda_tests.cu @@ -139,7 +139,7 @@ TEST(tALLPpmcVLReconstructor, CorrectInputExpectCorrectOutput) #ifdef MHD size_t const n_fields = 8; #else // not MHD - size_t const n_fields = 5; + size_t const n_fields = 5; #endif // MHD // Setup host grid. Fill host grid with random values and randomly assign maximum value diff --git a/src/reconstruction/ppmp_cuda.cu b/src/reconstruction/ppmp_cuda.cu index ae8da90cb..f84946437 100644 --- a/src/reconstruction/ppmp_cuda.cu +++ b/src/reconstruction/ppmp_cuda.cu @@ -166,7 +166,7 @@ __global__ void PPMP_cuda(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bou dge = dev_conserved[(n_fields - 1) * n_cells + id]; p_i = hydro_utilities::Get_Pressure_From_DE(E, E - E_kin, dge, gamma); #else - p_i = (dev_conserved[4 * n_cells + id] - 0.5 * d_i * (vx_i * vx_i + vy_i * vy_i + vz_i * vz_i)) * (gamma - 1.0); + p_i = (dev_conserved[4 * n_cells + id] - 0.5 * d_i * (vx_i * vx_i + vy_i * vy_i + vz_i * vz_i)) * (gamma - 1.0); #endif // PRESSURE_DE p_i = fmax(p_i, (Real)TINY_NUMBER); #ifdef DE diff --git a/src/reconstruction/reconstruction_tests.cu b/src/reconstruction/reconstruction_tests.cu index 6c2e19af7..74c0e6896 100644 --- a/src/reconstruction/reconstruction_tests.cu +++ b/src/reconstruction/reconstruction_tests.cu @@ -575,9 +575,13 @@ TEST(tALLReconstructionWriteData, CorrectInputExpectCorrectOutput) { // Set up test and mock up grid #ifdef MHD - reconstruction::Primitive interface{1, 2, 3, 4, 5, 6, 7, 8}; + reconstruction::Primitive interface { + 1, 2, 3, 4, 5, 6, 7, 8 + }; #else // MHD - reconstruction::Primitive interface{6, 7, 8, 9, 10}; + reconstruction::Primitive interface { + 6, 7, 8, 9, 10 + }; #endif // MHD size_t const nx = 3, ny = 3, nz = 3; size_t const n_cells = nx * ny * nz; diff --git a/src/system_tests/hydro_system_tests.cpp b/src/system_tests/hydro_system_tests.cpp index 6cffe9c21..18b2994bb 100644 --- a/src/system_tests/hydro_system_tests.cpp +++ b/src/system_tests/hydro_system_tests.cpp @@ -56,8 +56,8 @@ TEST_P(tHYDROtMHDSYSTEMSodShockTubeParameterizedMpi, CorrectInputExpectCorrectOu double const maxAllowedL1Error = 7.0E-3; double const maxAllowedError = 4.6E-2; #else - double const maxAllowedL1Error = 9.4E-5; - double const maxAllowedError = 6.4E-4; + double const maxAllowedL1Error = 9.4E-5; + double const maxAllowedError = 6.4E-4; #endif // MHD sodTest.numMpiRanks = GetParam(); From d737c4b407d723ed0da78537dfb84d80fb89c469 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Tue, 23 Jan 2024 16:24:40 -0500 Subject: [PATCH 23/40] run clang format --- src/analysis/feedback_analysis.cpp | 2 +- src/chemistry_gpu/chemistry_functions.cpp | 2 +- src/cooling_grackle/cool_grackle.cpp | 2 +- src/gravity/gravity_functions.cpp | 8 ++++---- src/gravity/gravity_functions_gpu.cu | 4 ++-- src/grid/grid3D.cpp | 8 ++++---- src/grid/initial_conditions.cpp | 20 ++++++++++---------- src/integrators/VL_3D_cuda.cu | 2 +- src/io/io.cpp | 4 ++-- src/particles/io_particles.cpp | 14 +++++++------- src/particles/particles_3D.cpp | 14 +++++++------- src/particles/particles_boundaries_cpu.cpp | 4 ++-- src/reconstruction/plmp_cuda.cu | 2 +- src/reconstruction/ppmc_cuda_tests.cu | 2 +- src/reconstruction/ppmp_cuda.cu | 2 +- src/reconstruction/reconstruction_tests.cu | 8 ++------ src/system_tests/hydro_system_tests.cpp | 4 ++-- 17 files changed, 49 insertions(+), 53 deletions(-) diff --git a/src/analysis/feedback_analysis.cpp b/src/analysis/feedback_analysis.cpp index 5fe0e7543..3dab7b6da 100644 --- a/src/analysis/feedback_analysis.cpp +++ b/src/analysis/feedback_analysis.cpp @@ -87,7 +87,7 @@ void FeedbackAnalysis::Compute_Gas_Velocity_Dispersion(Grid3D& G) #ifdef MPI_CHOLLA MPI_Allreduce(&partial_mass, &total_mass, 1, MPI_CHREAL, MPI_SUM, world); #else - total_mass = partial_mass; + total_mass = partial_mass; #endif for (k = G.H.n_ghost; k < G.H.nz - G.H.n_ghost; k++) { diff --git a/src/chemistry_gpu/chemistry_functions.cpp b/src/chemistry_gpu/chemistry_functions.cpp index 65e3af691..7999a6d55 100644 --- a/src/chemistry_gpu/chemistry_functions.cpp +++ b/src/chemistry_gpu/chemistry_functions.cpp @@ -228,7 +228,7 @@ void Grid3D::Update_Chemistry() #ifdef COSMOLOGY Chem.H.current_z = Cosmo.current_z; #else - Chem.H.current_z = 0; + Chem.H.current_z = 0; #endif Do_Chemistry_Update(C.device, H.nx, H.ny, H.nz, H.n_ghost, H.n_fields, H.dt, Chem.H); diff --git a/src/cooling_grackle/cool_grackle.cpp b/src/cooling_grackle/cool_grackle.cpp index c5f2a8078..a7f5c36cb 100644 --- a/src/cooling_grackle/cool_grackle.cpp +++ b/src/cooling_grackle/cool_grackle.cpp @@ -89,7 +89,7 @@ void Cool_GK::Initialize(struct Parameters *P, Cosmology &Cosmo) data->metal_cooling = 1; // metal cooling off #else chprintf("WARNING: Metal Cooling is Off. \n"); - data->metal_cooling = 0; // metal cooling off + data->metal_cooling = 0; // metal cooling off #endif #ifdef PARALLEL_OMP diff --git a/src/gravity/gravity_functions.cpp b/src/gravity/gravity_functions.cpp index 6c7a6dde7..b92d06564 100644 --- a/src/gravity/gravity_functions.cpp +++ b/src/gravity/gravity_functions.cpp @@ -137,7 +137,7 @@ void Grid3D::set_dt_Gravity() dt_particles = Calc_Particles_dt(); dt_particles = fmin(dt_particles, Particles.max_dt); #ifdef ONLY_PARTICLES - dt_min = dt_particles; + dt_min = dt_particles; chprintf(" dt_particles: %f \n", dt_particles); #else chprintf(" dt_hydro: %f dt_particles: %f \n", dt_hydro, dt_particles); @@ -211,7 +211,7 @@ Real Grav3D::Get_Average_Density() #ifdef MPI_CHOLLA dens_avrg_all = ReduceRealAvg(dens_mean); #else - dens_avrg_all = dens_mean; + dens_avrg_all = dens_mean; #endif dens_avrg = dens_avrg_all; @@ -530,8 +530,8 @@ void Grid3D::Compute_Gravitational_Potential(struct Parameters *P) input_density = Grav.F.density_d; output_potential = Grav.F.potential_d; #else - input_density = Grav.F.density_h; - output_potential = Grav.F.potential_h; + input_density = Grav.F.density_h; + output_potential = Grav.F.potential_h; #endif #ifdef SOR diff --git a/src/gravity/gravity_functions_gpu.cu b/src/gravity/gravity_functions_gpu.cu index 15de64a95..b92d19084 100644 --- a/src/gravity/gravity_functions_gpu.cu +++ b/src/gravity/gravity_functions_gpu.cu @@ -127,7 +127,7 @@ void Grid3D::Copy_Hydro_Density_to_Gravity_GPU() #ifdef COSMOLOGY cosmo_rho_0_gas = Cosmo.rho_0_gas; #else - cosmo_rho_0_gas = 1.0; + cosmo_rho_0_gas = 1.0; #endif // Copy the density from the device array to the Poisson input density array @@ -261,7 +261,7 @@ void Grid3D::Extrapolate_Grav_Potential_GPU() #ifdef COSMOLOGY cosmo_factor = Cosmo.current_a * Cosmo.current_a / Cosmo.phi_0_gas; #else - cosmo_factor = 1.0; + cosmo_factor = 1.0; #endif // set values for GPU kernels diff --git a/src/grid/grid3D.cpp b/src/grid/grid3D.cpp index 93499fea0..dd65da756 100644 --- a/src/grid/grid3D.cpp +++ b/src/grid/grid3D.cpp @@ -267,13 +267,13 @@ void Grid3D::Initialize(struct Parameters *P) #ifdef DENSITY_FLOOR H.density_floor = P->density_floor; #else - H.density_floor = 0.0; + H.density_floor = 0.0; #endif #ifdef SCALAR_FLOOR H.scalar_floor = P->scalar_floor; #else - H.scalar_floor = 0.0; + H.scalar_floor = 0.0; #endif #ifdef COSMOLOGY @@ -345,8 +345,8 @@ void Grid3D::AllocateMemory(void) GPU_Error_Check(cudaHostAlloc(&C.Grav_potential, H.n_cells * sizeof(Real), cudaHostAllocDefault)); GPU_Error_Check(cudaMalloc((void **)&C.d_Grav_potential, H.n_cells * sizeof(Real))); #else - C.Grav_potential = NULL; - C.d_Grav_potential = NULL; + C.Grav_potential = NULL; + C.d_Grav_potential = NULL; #endif #ifdef CHEMISTRY_GPU diff --git a/src/grid/initial_conditions.cpp b/src/grid/initial_conditions.cpp index 30e3eb459..af558be8f 100644 --- a/src/grid/initial_conditions.cpp +++ b/src/grid/initial_conditions.cpp @@ -1482,18 +1482,18 @@ void Grid3D::Zeldovich_Pancake(struct Parameters P) Real H0, h, Omega_M, rho_0, G, z_zeldovich, z_init, x_center, T_init, k_x; chprintf("Setting Zeldovich Pancake initial conditions...\n"); - H0 = P.H0; - h = H0 / 100; + H0 = P.H0; + h = H0 / 100; Omega_M = P.Omega_M; chprintf(" h = %f \n", h); chprintf(" Omega_M = %f \n", Omega_M); H0 /= 1000; //[km/s / kpc] - G = G_COSMO; - rho_0 = 3 * H0 * H0 / (8 * M_PI * G) * Omega_M / h / h; + G = G_COSMO; + rho_0 = 3 * H0 * H0 / (8 * M_PI * G) * Omega_M / h / h; z_zeldovich = 1; - z_init = P.Init_redshift; + z_init = P.Init_redshift; chprintf(" rho_0 = %f \n", rho_0); chprintf(" z_init = %f \n", z_init); chprintf(" z_zeldovich = %f \n", z_zeldovich); @@ -1553,17 +1553,17 @@ void Grid3D::Zeldovich_Pancake(struct Parameters P) index = (int(x_pos / H.dx) + 0) % 256; // index = ( index + 16 ) % 256; dens = ics_values[0 * nPoints + index]; - vel = ics_values[1 * nPoints + index]; - E = ics_values[2 * nPoints + index]; - U = ics_values[3 * nPoints + index]; + vel = ics_values[1 * nPoints + index]; + E = ics_values[2 * nPoints + index]; + U = ics_values[3 * nPoints + index]; // // // chprintf( "%f \n", vel ); - C.density[id] = dens; + C.density[id] = dens; C.momentum_x[id] = dens * vel; C.momentum_y[id] = 0; C.momentum_z[id] = 0; - C.Energy[id] = E; + C.Energy[id] = E; #ifdef DE C.GasEnergy[id] = U; diff --git a/src/integrators/VL_3D_cuda.cu b/src/integrators/VL_3D_cuda.cu index 2b4162ffc..1118c1419 100644 --- a/src/integrators/VL_3D_cuda.cu +++ b/src/integrators/VL_3D_cuda.cu @@ -123,7 +123,7 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int #if defined(GRAVITY) dev_grav_potential = d_grav_potential; #else // not GRAVITY - dev_grav_potential = NULL; + dev_grav_potential = NULL; #endif // GRAVITY // If memory is single allocated: memory_allocated becomes true and diff --git a/src/io/io.cpp b/src/io/io.cpp index b243440e4..9959267de 100644 --- a/src/io/io.cpp +++ b/src/io/io.cpp @@ -1400,12 +1400,12 @@ void Grid3D::Write_Grid_HDF5(hid_t file_id) #ifdef OUTPUT_METALS output_metals = true; #else // not OUTPUT_METALS - output_metals = false; + output_metals = false; #endif // OUTPUT_METALS #ifdef OUTPUT_ELECTRONS output_electrons = true; #else // not OUTPUT_ELECTRONS - output_electrons = false; + output_electrons = false; #endif // OUTPUT_ELECTRONS #ifdef OUTPUT_FULL_IONIZATION output_full_ionization = true; diff --git a/src/particles/io_particles.cpp b/src/particles/io_particles.cpp index e6da774ed..e986c5287 100644 --- a/src/particles/io_particles.cpp +++ b/src/particles/io_particles.cpp @@ -445,12 +445,12 @@ void Particles3D::Load_Particles_Data_HDF5(hid_t file_id, int nfile, struct Para Real vy_max_g = vy_max; Real vz_max_g = vz_max; - Real px_min_g = px_min; - Real py_min_g = py_min; - Real pz_min_g = pz_min; - Real vx_min_g = vx_min; - Real vy_min_g = vy_min; - Real vz_min_g = vz_min; + Real px_min_g = px_min; + Real py_min_g = py_min; + Real pz_min_g = pz_min; + Real vx_min_g = vx_min; + Real vy_min_g = vy_min; + Real vz_min_g = vz_min; #endif // MPI_CHOLLA // Print initial Statistics @@ -563,7 +563,7 @@ void Grid3D::Write_Particles_Data_HDF5(hid_t file_id) #ifdef MPI_CHOLLA N_particles_total = ReducePartIntSum(Particles.n_local); #else - N_particles_total = Particles.n_local; + N_particles_total = Particles.n_local; #endif // Print the total particles when saving the particles data diff --git a/src/particles/particles_3D.cpp b/src/particles/particles_3D.cpp index 87a2be8e5..6417e4136 100644 --- a/src/particles/particles_3D.cpp +++ b/src/particles/particles_3D.cpp @@ -157,12 +157,12 @@ void Particles3D::Initialize(struct Parameters *P, Grav3D &Grav, Real xbound, Re G.boundary_type_z0 = P->zlg_bcnd; G.boundary_type_z1 = P->zug_bcnd; #else - G.boundary_type_x0 = P->xl_bcnd; - G.boundary_type_x1 = P->xu_bcnd; - G.boundary_type_y0 = P->yl_bcnd; - G.boundary_type_y1 = P->yu_bcnd; - G.boundary_type_z0 = P->zl_bcnd; - G.boundary_type_z1 = P->zu_bcnd; + G.boundary_type_x0 = P->xl_bcnd; + G.boundary_type_x1 = P->xu_bcnd; + G.boundary_type_y0 = P->yl_bcnd; + G.boundary_type_y1 = P->yu_bcnd; + G.boundary_type_z0 = P->zl_bcnd; + G.boundary_type_z1 = P->zu_bcnd; #endif #ifdef PARTICLES_GPU @@ -211,7 +211,7 @@ void Particles3D::Initialize(struct Parameters *P, Grav3D &Grav, Real xbound, Re #ifdef MPI_CHOLLA n_total_initial = ReducePartIntSum(n_local); #else - n_total_initial = n_local; + n_total_initial = n_local; #endif chprintf("Particles Initialized: \n n_local: %lu \n", n_local); diff --git a/src/particles/particles_boundaries_cpu.cpp b/src/particles/particles_boundaries_cpu.cpp index 772153534..27470befe 100644 --- a/src/particles/particles_boundaries_cpu.cpp +++ b/src/particles/particles_boundaries_cpu.cpp @@ -433,13 +433,13 @@ void Particles3D::Unload_Particles_from_Buffer_CPU(int direction, int side, Real offset_extra += 1; pId = recv_buffer[offset_extra]; #else - pId = 0; + pId = 0; #endif #ifdef PARTICLE_AGE offset_extra += 1; pAge = recv_buffer[offset_extra]; #else - pAge = 0.0; + pAge = 0.0; #endif offset_buff += N_DATA_PER_PARTICLE_TRANSFER; diff --git a/src/reconstruction/plmp_cuda.cu b/src/reconstruction/plmp_cuda.cu index f69bbdc4b..a000da4da 100644 --- a/src/reconstruction/plmp_cuda.cu +++ b/src/reconstruction/plmp_cuda.cu @@ -120,7 +120,7 @@ __global__ void PLMP_cuda(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bou dge = dev_conserved[(n_fields - 1) * n_cells + id]; p_i = hydro_utilities::Get_Pressure_From_DE(E, E - E_kin, dge, gamma); #else - p_i = (dev_conserved[4 * n_cells + id] - 0.5 * d_i * (vx_i * vx_i + vy_i * vy_i + vz_i * vz_i)) * (gamma - 1.0); + p_i = (dev_conserved[4 * n_cells + id] - 0.5 * d_i * (vx_i * vx_i + vy_i * vy_i + vz_i * vz_i)) * (gamma - 1.0); #endif // PRESSURE_DE p_i = fmax(p_i, (Real)TINY_NUMBER); #ifdef SCALAR diff --git a/src/reconstruction/ppmc_cuda_tests.cu b/src/reconstruction/ppmc_cuda_tests.cu index c1319ea58..9e9b11140 100644 --- a/src/reconstruction/ppmc_cuda_tests.cu +++ b/src/reconstruction/ppmc_cuda_tests.cu @@ -139,7 +139,7 @@ TEST(tALLPpmcVLReconstructor, CorrectInputExpectCorrectOutput) #ifdef MHD size_t const n_fields = 8; #else // not MHD - size_t const n_fields = 5; + size_t const n_fields = 5; #endif // MHD // Setup host grid. Fill host grid with random values and randomly assign maximum value diff --git a/src/reconstruction/ppmp_cuda.cu b/src/reconstruction/ppmp_cuda.cu index f84946437..ae8da90cb 100644 --- a/src/reconstruction/ppmp_cuda.cu +++ b/src/reconstruction/ppmp_cuda.cu @@ -166,7 +166,7 @@ __global__ void PPMP_cuda(Real *dev_conserved, Real *dev_bounds_L, Real *dev_bou dge = dev_conserved[(n_fields - 1) * n_cells + id]; p_i = hydro_utilities::Get_Pressure_From_DE(E, E - E_kin, dge, gamma); #else - p_i = (dev_conserved[4 * n_cells + id] - 0.5 * d_i * (vx_i * vx_i + vy_i * vy_i + vz_i * vz_i)) * (gamma - 1.0); + p_i = (dev_conserved[4 * n_cells + id] - 0.5 * d_i * (vx_i * vx_i + vy_i * vy_i + vz_i * vz_i)) * (gamma - 1.0); #endif // PRESSURE_DE p_i = fmax(p_i, (Real)TINY_NUMBER); #ifdef DE diff --git a/src/reconstruction/reconstruction_tests.cu b/src/reconstruction/reconstruction_tests.cu index 74c0e6896..6c2e19af7 100644 --- a/src/reconstruction/reconstruction_tests.cu +++ b/src/reconstruction/reconstruction_tests.cu @@ -575,13 +575,9 @@ TEST(tALLReconstructionWriteData, CorrectInputExpectCorrectOutput) { // Set up test and mock up grid #ifdef MHD - reconstruction::Primitive interface { - 1, 2, 3, 4, 5, 6, 7, 8 - }; + reconstruction::Primitive interface{1, 2, 3, 4, 5, 6, 7, 8}; #else // MHD - reconstruction::Primitive interface { - 6, 7, 8, 9, 10 - }; + reconstruction::Primitive interface{6, 7, 8, 9, 10}; #endif // MHD size_t const nx = 3, ny = 3, nz = 3; size_t const n_cells = nx * ny * nz; diff --git a/src/system_tests/hydro_system_tests.cpp b/src/system_tests/hydro_system_tests.cpp index 18b2994bb..6cffe9c21 100644 --- a/src/system_tests/hydro_system_tests.cpp +++ b/src/system_tests/hydro_system_tests.cpp @@ -56,8 +56,8 @@ TEST_P(tHYDROtMHDSYSTEMSodShockTubeParameterizedMpi, CorrectInputExpectCorrectOu double const maxAllowedL1Error = 7.0E-3; double const maxAllowedError = 4.6E-2; #else - double const maxAllowedL1Error = 9.4E-5; - double const maxAllowedError = 6.4E-4; + double const maxAllowedL1Error = 9.4E-5; + double const maxAllowedError = 6.4E-4; #endif // MHD sodTest.numMpiRanks = GetParam(); From d292c46cbfd2fe00a3eb12603bc2c0d3c25bd1ef Mon Sep 17 00:00:00 2001 From: Helena Richie Date: Tue, 23 Jan 2024 18:08:43 -0500 Subject: [PATCH 24/40] undo accidental change --- src/grid/grid3D.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/grid/grid3D.cpp b/src/grid/grid3D.cpp index 1cb8d9f56..4c3ddd1ab 100644 --- a/src/grid/grid3D.cpp +++ b/src/grid/grid3D.cpp @@ -157,7 +157,7 @@ void Grid3D::Initialize(struct Parameters *P) C_cfl = 0.3; #ifdef AVERAGE_SLOW_CELLS - H.min_dt_slow = 1e-5; // Initialize the minumum dt to a tiny number + H.min_dt_slow = 1e-100; // Initialize the minumum dt to a tiny number #endif // AVERAGE_SLOW_CELLS #ifndef MPI_CHOLLA From da2d264db7b148366e7cb31e3cb83885ccc94161 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Tue, 23 Jan 2024 18:38:35 -0500 Subject: [PATCH 25/40] undo accidental removal of density floor in half-step --- cholla-tests-data | 2 +- src/integrators/VL_3D_cuda.cu | 15 +++++++++++++++ 2 files changed, 16 insertions(+), 1 deletion(-) diff --git a/cholla-tests-data b/cholla-tests-data index 71eb66d63..dcd73ff52 160000 --- a/cholla-tests-data +++ b/cholla-tests-data @@ -1 +1 @@ -Subproject commit 71eb66d63622ac15c0844ae96ec9034cd5e4f4d3 +Subproject commit dcd73ff52b9027627b247c6d888bcdb56840c85e diff --git a/src/integrators/VL_3D_cuda.cu b/src/integrators/VL_3D_cuda.cu index 1118c1419..b1bc118b7 100644 --- a/src/integrators/VL_3D_cuda.cu +++ b/src/integrators/VL_3D_cuda.cu @@ -455,6 +455,21 @@ __global__ void Update_Conserved_Variables_3D_half(Real *dev_conserved, Real *de dtodz * (dev_F_z[(n_fields - 1) * n_cells + kmo] - dev_F_z[(n_fields - 1) * n_cells + id]) + 0.5 * P * (dtodx * (vx_imo - vx_ipo) + dtody * (vy_jmo - vy_jpo) + dtodz * (vz_kmo - vz_kpo)); #endif // DE + #ifdef DENSITY_FLOOR + if (dev_conserved_half[id] < density_floor) { + dens_0 = dev_conserved_half[id]; + printf("###Thread density change %f -> %f \n", dens_0, density_floor); + dev_conserved_half[id] = density_floor; + // Scale the conserved values to the new density + dev_conserved_half[1 * n_cells + id] *= (density_floor / dens_0); + dev_conserved_half[2 * n_cells + id] *= (density_floor / dens_0); + dev_conserved_half[3 * n_cells + id] *= (density_floor / dens_0); + dev_conserved_half[4 * n_cells + id] *= (density_floor / dens_0); + #ifdef DE + dev_conserved_half[(n_fields - 1) * n_cells + id] *= (density_floor / dens_0); + #endif // DE + } + #endif // DENSITY_FLOOR } } From ef470301b2091ab7c5159e9d3908acb21c7bf687 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Tue, 23 Jan 2024 18:49:16 -0500 Subject: [PATCH 26/40] fix undeclared variable --- src/integrators/VL_3D_cuda.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/integrators/VL_3D_cuda.cu b/src/integrators/VL_3D_cuda.cu index b1bc118b7..0657a946f 100644 --- a/src/integrators/VL_3D_cuda.cu +++ b/src/integrators/VL_3D_cuda.cu @@ -457,7 +457,7 @@ __global__ void Update_Conserved_Variables_3D_half(Real *dev_conserved, Real *de #endif // DE #ifdef DENSITY_FLOOR if (dev_conserved_half[id] < density_floor) { - dens_0 = dev_conserved_half[id]; + Real dens_0 = dev_conserved_half[id]; printf("###Thread density change %f -> %f \n", dens_0, density_floor); dev_conserved_half[id] = density_floor; // Scale the conserved values to the new density From 2a5d7ed34b1c9d268827b1014af407d7662e5bd6 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Tue, 23 Jan 2024 19:41:30 -0500 Subject: [PATCH 27/40] remove debugging print statement --- src/dust/dust_cuda.cu | 1 - 1 file changed, 1 deletion(-) diff --git a/src/dust/dust_cuda.cu b/src/dust/dust_cuda.cu index 28f764900..8376a2a1b 100644 --- a/src/dust/dust_cuda.cu +++ b/src/dust/dust_cuda.cu @@ -139,7 +139,6 @@ __device__ __host__ Real Calc_Sputtering_Timescale(Real number_density, Real tem number_density /= (6e-4); // gas number density in units of 10^-27 g/cm^3 // sputtering timescale, s - printf("%e\n", grain_radius); Real tau_sp = A * (a / number_density) * (pow(temperature_0 / temperature, omega) + 1); return tau_sp; From db0f78865d64ca77a0a9d05cc3043840080d02be Mon Sep 17 00:00:00 2001 From: helenarichie Date: Tue, 23 Jan 2024 19:50:35 -0500 Subject: [PATCH 28/40] add host wrapper functions for temp and scalar floor kernels and move calls to grid3d::update_hydro_grid --- src/grid/grid3D.cpp | 38 ++++++++++++++++++------------- src/hydro/hydro_cuda.cu | 37 ++++++++++++++++++++++++++---- src/hydro/hydro_cuda.h | 8 +++++-- src/hydro/hydro_cuda_tests.cu | 6 ++--- src/integrators/VL_3D_cuda.cu | 18 ++------------- src/integrators/VL_3D_cuda.h | 4 ++-- src/integrators/simple_3D_cuda.cu | 18 ++------------- src/integrators/simple_3D_cuda.h | 4 ++-- 8 files changed, 71 insertions(+), 62 deletions(-) diff --git a/src/grid/grid3D.cpp b/src/grid/grid3D.cpp index 4c3ddd1ab..2e548e1f6 100644 --- a/src/grid/grid3D.cpp +++ b/src/grid/grid3D.cpp @@ -430,16 +430,6 @@ void Grid3D::Execute_Hydro_Integrator(void) z_off = nz_local_start; #endif - // Set the lower limit for density and temperature (Internal Energy) - Real U_floor, density_floor; - density_floor = H.density_floor; - // Minimum of internal energy from minumum of temperature - U_floor = H.temperature_floor * KB / (gama - 1) / MP / SP_ENERGY_UNIT; -#ifdef COSMOLOGY - U_floor = H.temperature_floor / (gama - 1) / MP * KB * 1e-10; // ( km/s )^2 - U_floor /= Cosmo.v_0_gas * Cosmo.v_0_gas / Cosmo.current_a / Cosmo.current_a; -#endif - #ifdef CPU_TIME Timer.Hydro_Integrator.Start(); #endif // CPU_TIME @@ -472,13 +462,12 @@ void Grid3D::Execute_Hydro_Integrator(void) #ifdef CUDA #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, density_floor, U_floor, - C.Grav_potential, H.scalar_floor); + H.dz, H.xbound, H.ybound, H.zbound, H.dt, H.n_fields, H.custom_grav, H.density_floor, + C.Grav_potential); #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, - H.dz, H.xbound, H.ybound, H.zbound, H.dt, H.n_fields, H.custom_grav, density_floor, - U_floor, C.Grav_potential, H.scalar_floor); + H.dz, H.xbound, H.ybound, H.zbound, H.dt, H.n_fields, H.custom_grav, H.density_floor, C.Grav_potential); #endif // SIMPLE #endif } else { @@ -512,10 +501,27 @@ Real Grid3D::Update_Hydro_Grid() Execute_Hydro_Integrator(); - // == Perform chemistry/cooling (there are a few different cases) == - #ifdef CUDA + #ifdef TEMPERATURE_FLOOR + // Set the lower limit temperature (Internal Energy) + Real U_floor; + // Minimum of internal energy from minumum of temperature + U_floor = H.temperature_floor * KB / (gama - 1) / MP / SP_ENERGY_UNIT; + #ifdef COSMOLOGY + U_floor = H.temperature_floor / (gama - 1) / MP * KB * 1e-10; // ( km/s )^2 + U_floor /= Cosmo.v_0_gas * Cosmo.v_0_gas / Cosmo.current_a / Cosmo.current_a; + #endif + Apply_Temperature_Floor(C.device, H.nx, H.ny, H.nz, H.n_ghost, H.n_fields, U_floor); + #endif // TEMPERATURE_FLOOR + + #ifdef SCALAR_FLOOR + #ifdef DUST + Apply_Scalar_Floor(C.device, H.nx, H.ny, H.nz, H.n_ghost, grid_enum::dust_density, H.scalar_floor); + #endif + #endif // SCALAR_FLOOR + + // == Perform chemistry/cooling (there are a few different cases) == #ifdef COOLING_GPU #ifdef CPU_TIME Timer.Cooling_GPU.Start(); diff --git a/src/hydro/hydro_cuda.cu b/src/hydro/hydro_cuda.cu index 125e851b4..484023353 100644 --- a/src/hydro/hydro_cuda.cu +++ b/src/hydro/hydro_cuda.cu @@ -1101,8 +1101,22 @@ __global__ void Sync_Energies_3D(Real *dev_conserved, int nx, int ny, int nz, in #endif // DE -__global__ void Apply_Temperature_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, - Real U_floor) + +void Apply_Temperature_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real U_floor) +{ + // set values for GPU kernels + int n_cells = nx * ny * nz; + int ngrid = (n_cells + TPB - 1) / TPB; + // number of blocks per 1D grid + dim3 dim1dGrid(ngrid, 1, 1); + // number of threads per 1D block + dim3 dim1dBlock(TPB, 1, 1); + + hipLaunchKernelGGL(Temperature_Floor_Kernel, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, n_fields, U_floor); +} + +__global__ void Temperature_Floor_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, + Real U_floor) { int id, xid, yid, zid, n_cells; Real d, d_inv, vx, vy, vz, E, Ekin, U; @@ -1254,8 +1268,21 @@ __device__ void Average_Cell_All_Fields(int i, int j, int k, int nx, int ny, int printf("%3d %3d %3d FC: d: %e E:%e P:%e vx:%e vy:%e vz:%e\n", i, j, k, d, E, P, vx_av, vy_av, vz_av); } -__global__ void Apply_Scalar_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, - Real scalar_floor) +void Apply_Scalar_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, Real scalar_floor) +{ + // set values for GPU kernels + int n_cells = nx * ny * nz; + int ngrid = (n_cells + TPB - 1) / TPB; + // number of blocks per 1D grid + dim3 dim1dGrid(ngrid, 1, 1); + // number of threads per 1D block + dim3 dim1dBlock(TPB, 1, 1); + + hipLaunchKernelGGL(Scalar_Floor_Kernel, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, field_num, scalar_floor); +} + +__global__ void Scalar_Floor_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, + Real scalar_floor) { int id, xid, yid, zid, n_cells; Real scalar; // variable to store the value of the scalar before a floor is applied @@ -1273,7 +1300,7 @@ __global__ void Apply_Scalar_Floor(Real *dev_conserved, int nx, int ny, int nz, scalar = dev_conserved[id + n_cells * field_num]; if (scalar < scalar_floor) { - printf("###Thread scalar change %f -> %f \n", scalar, scalar_floor); + // printf("###Thread scalar change %f -> %f \n", scalar, scalar_floor); dev_conserved[id + n_cells * field_num] = scalar_floor; } } diff --git a/src/hydro/hydro_cuda.h b/src/hydro/hydro_cuda.h index 8fcfbba05..ea64b2164 100644 --- a/src/hydro/hydro_cuda.h +++ b/src/hydro/hydro_cuda.h @@ -85,10 +85,14 @@ __global__ void Average_Slow_Cells_3D(Real *dev_conserved, int nx, int ny, int n Real dy, Real dz, Real gamma, Real max_dti_slow); #endif -__global__ void Apply_Temperature_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, +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, Real U_floor); -__global__ void Apply_Scalar_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, +void Apply_Scalar_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, Real scalar_floor); + +__global__ void Scalar_Floor_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, Real scalar_floor); __global__ void Partial_Update_Advected_Internal_Energy_1D(Real *dev_conserved, Real *Q_Lx, Real *Q_Rx, int nx, diff --git a/src/hydro/hydro_cuda_tests.cu b/src/hydro/hydro_cuda_tests.cu index 0796a3064..ade162926 100644 --- a/src/hydro/hydro_cuda_tests.cu +++ b/src/hydro/hydro_cuda_tests.cu @@ -174,21 +174,21 @@ TEST(tHYDROScalarFloor, CorrectInputExpectCorrectOutput) // Case where scalar is below the floor host_conserved.at(field_num) = 0.0; // scalar dev_conserved.cpyHostToDevice(host_conserved); - hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved.data(), nx, ny, nz, n_ghost, + hipLaunchKernelGGL(Scalar_Floor_Kernel, dim1dGrid, dim1dBlock, 0, 0, dev_conserved.data(), nx, ny, nz, n_ghost, field_num, scalar_floor); testing_utilities::Check_Results(scalar_floor, dev_conserved.at(field_num), "below floor"); // Case where scalar is above the floor host_conserved.at(field_num) = 2.0; // scalar dev_conserved.cpyHostToDevice(host_conserved); - hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved.data(), nx, ny, nz, n_ghost, + hipLaunchKernelGGL(Scalar_Floor_Kernel, dim1dGrid, dim1dBlock, 0, 0, dev_conserved.data(), nx, ny, nz, n_ghost, field_num, scalar_floor); testing_utilities::Check_Results(host_conserved.at(field_num), dev_conserved.at(field_num), "above floor"); // Case where scalar is at the floor host_conserved.at(field_num) = 1.0; // scalar dev_conserved.cpyHostToDevice(host_conserved); - hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved.data(), nx, ny, nz, n_ghost, + hipLaunchKernelGGL(Scalar_Floor_Kernel, dim1dGrid, dim1dBlock, 0, 0, dev_conserved.data(), nx, ny, nz, n_ghost, field_num, scalar_floor); testing_utilities::Check_Results(host_conserved.at(field_num), dev_conserved.at(field_num), "at floor"); } diff --git a/src/integrators/VL_3D_cuda.cu b/src/integrators/VL_3D_cuda.cu index 0657a946f..c0250983f 100644 --- a/src/integrators/VL_3D_cuda.cu +++ b/src/integrators/VL_3D_cuda.cu @@ -37,8 +37,8 @@ __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 U_floor, - Real *host_grav_potential, Real scalar_floor) + Real dt, int n_fields, int custom_grav, Real density_floor, + Real *host_grav_potential) { // Here, *dev_conserved contains the entire // set of conserved variables on the grid @@ -321,20 +321,6 @@ void VL_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, int GPU_Error_Check(); #endif // DE - #ifdef TEMPERATURE_FLOOR - hipLaunchKernelGGL(Apply_Temperature_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, n_fields, - U_floor); - GPU_Error_Check(); - #endif // TEMPERATURE_FLOOR - - #ifdef SCALAR_FLOOR - #ifdef DUST - hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, - grid_enum::dust_density, scalar_floor); - GPU_Error_Check(); - #endif - #endif // SCALAR_FLOOR - return; } diff --git a/src/integrators/VL_3D_cuda.h b/src/integrators/VL_3D_cuda.h index 4104493bc..3310f5fe5 100644 --- a/src/integrators/VL_3D_cuda.h +++ b/src/integrators/VL_3D_cuda.h @@ -10,8 +10,8 @@ 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 U_floor, - Real *host_grav_potential, Real scalar_floor); + Real dt, int n_fields, int custom_grav, Real density_floor, + Real *host_grav_potential); void Free_Memory_VL_3D(); diff --git a/src/integrators/simple_3D_cuda.cu b/src/integrators/simple_3D_cuda.cu index b68f0a351..3a1ce77b1 100644 --- a/src/integrators/simple_3D_cuda.cu +++ b/src/integrators/simple_3D_cuda.cu @@ -26,8 +26,8 @@ void Simple_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 U_floor, - Real *host_grav_potential, Real scalar_floor) + Real dt, int n_fields, int custom_grav, Real density_floor, + Real *host_grav_potential) { // Here, *dev_conserved contains the entire // set of conserved variables on the grid @@ -180,20 +180,6 @@ void Simple_Algorithm_3D_CUDA(Real *d_conserved, Real *d_grav_potential, int nx, GPU_Error_Check(); #endif - #ifdef TEMPERATURE_FLOOR - hipLaunchKernelGGL(Apply_Temperature_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, n_fields, - U_floor); - GPU_Error_Check(); - #endif // TEMPERATURE_FLOOR - - #ifdef SCALAR_FLOOR - #ifdef DUST - hipLaunchKernelGGL(Apply_Scalar_Floor, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, - grid_enum::dust_density, scalar_floor); - CudaCheckError(); - #endif DUST - #endif // SCALAR_FLOOR - return; } diff --git a/src/integrators/simple_3D_cuda.h b/src/integrators/simple_3D_cuda.h index e2cea247e..57d211ed1 100644 --- a/src/integrators/simple_3D_cuda.h +++ b/src/integrators/simple_3D_cuda.h @@ -11,8 +11,8 @@ void Simple_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 U_floor, - Real *host_grav_potential, Real scalar_floor); + Real dt, int n_fields, int custom_grav, Real density_floor, + Real *host_grav_potential); void Free_Memory_Simple_3D(); From 71df1156e81ac3c51705448e74fb6e2ec77a3ede Mon Sep 17 00:00:00 2001 From: helenarichie Date: Tue, 23 Jan 2024 19:51:16 -0500 Subject: [PATCH 29/40] run clang format' --- src/grid/grid3D.cpp | 11 ++++++----- src/hydro/hydro_cuda.cu | 11 ++++++----- src/hydro/hydro_cuda.h | 4 ++-- src/integrators/VL_3D_cuda.cu | 3 +-- src/integrators/VL_3D_cuda.h | 3 +-- src/integrators/simple_3D_cuda.cu | 3 +-- src/integrators/simple_3D_cuda.h | 3 +-- 7 files changed, 18 insertions(+), 20 deletions(-) diff --git a/src/grid/grid3D.cpp b/src/grid/grid3D.cpp index 2e548e1f6..d25189d89 100644 --- a/src/grid/grid3D.cpp +++ b/src/grid/grid3D.cpp @@ -158,7 +158,7 @@ void Grid3D::Initialize(struct Parameters *P) #ifdef AVERAGE_SLOW_CELLS H.min_dt_slow = 1e-100; // Initialize the minumum dt to a tiny number -#endif // AVERAGE_SLOW_CELLS +#endif // AVERAGE_SLOW_CELLS #ifndef MPI_CHOLLA @@ -467,7 +467,8 @@ void Grid3D::Execute_Hydro_Integrator(void) #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, - H.dz, H.xbound, H.ybound, H.zbound, H.dt, H.n_fields, H.custom_grav, H.density_floor, C.Grav_potential); + H.dz, H.xbound, H.ybound, H.zbound, H.dt, H.n_fields, H.custom_grav, H.density_floor, + C.Grav_potential); #endif // SIMPLE #endif } else { @@ -508,16 +509,16 @@ Real Grid3D::Update_Hydro_Grid() Real U_floor; // Minimum of internal energy from minumum of temperature U_floor = H.temperature_floor * KB / (gama - 1) / MP / SP_ENERGY_UNIT; - #ifdef COSMOLOGY + #ifdef COSMOLOGY U_floor = H.temperature_floor / (gama - 1) / MP * KB * 1e-10; // ( km/s )^2 U_floor /= Cosmo.v_0_gas * Cosmo.v_0_gas / Cosmo.current_a / Cosmo.current_a; - #endif + #endif Apply_Temperature_Floor(C.device, H.nx, H.ny, H.nz, H.n_ghost, H.n_fields, U_floor); #endif // TEMPERATURE_FLOOR #ifdef SCALAR_FLOOR #ifdef DUST - Apply_Scalar_Floor(C.device, H.nx, H.ny, H.nz, H.n_ghost, grid_enum::dust_density, H.scalar_floor); + Apply_Scalar_Floor(C.device, H.nx, H.ny, H.nz, H.n_ghost, grid_enum::dust_density, H.scalar_floor); #endif #endif // SCALAR_FLOOR diff --git a/src/hydro/hydro_cuda.cu b/src/hydro/hydro_cuda.cu index 484023353..69c4c083d 100644 --- a/src/hydro/hydro_cuda.cu +++ b/src/hydro/hydro_cuda.cu @@ -1101,7 +1101,6 @@ __global__ void Sync_Energies_3D(Real *dev_conserved, int nx, int ny, int nz, in #endif // DE - void Apply_Temperature_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, Real U_floor) { // set values for GPU kernels @@ -1111,8 +1110,9 @@ void Apply_Temperature_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ dim3 dim1dGrid(ngrid, 1, 1); // number of threads per 1D block dim3 dim1dBlock(TPB, 1, 1); - - hipLaunchKernelGGL(Temperature_Floor_Kernel, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, n_fields, U_floor); + + hipLaunchKernelGGL(Temperature_Floor_Kernel, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, + n_fields, U_floor); } __global__ void Temperature_Floor_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n_fields, @@ -1277,8 +1277,9 @@ void Apply_Scalar_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost dim3 dim1dGrid(ngrid, 1, 1); // number of threads per 1D block dim3 dim1dBlock(TPB, 1, 1); - - hipLaunchKernelGGL(Scalar_Floor_Kernel, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, field_num, scalar_floor); + + hipLaunchKernelGGL(Scalar_Floor_Kernel, dim1dGrid, dim1dBlock, 0, 0, dev_conserved, nx, ny, nz, n_ghost, field_num, + scalar_floor); } __global__ void Scalar_Floor_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, diff --git a/src/hydro/hydro_cuda.h b/src/hydro/hydro_cuda.h index ea64b2164..c6d9ac96d 100644 --- a/src/hydro/hydro_cuda.h +++ b/src/hydro/hydro_cuda.h @@ -88,12 +88,12 @@ __global__ void Average_Slow_Cells_3D(Real *dev_conserved, int nx, int ny, int n 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, - Real U_floor); + Real U_floor); void Apply_Scalar_Floor(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, Real scalar_floor); __global__ void Scalar_Floor_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int field_num, - Real scalar_floor); + Real scalar_floor); __global__ void Partial_Update_Advected_Internal_Energy_1D(Real *dev_conserved, Real *Q_Lx, Real *Q_Rx, int nx, int n_ghost, Real dx, Real dt, Real gamma, int n_fields); diff --git a/src/integrators/VL_3D_cuda.cu b/src/integrators/VL_3D_cuda.cu index c0250983f..3d1c8eb11 100644 --- a/src/integrators/VL_3D_cuda.cu +++ b/src/integrators/VL_3D_cuda.cu @@ -37,8 +37,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) + Real dt, int n_fields, int custom_grav, Real density_floor, Real *host_grav_potential) { // Here, *dev_conserved contains the entire // set of conserved variables on the grid diff --git a/src/integrators/VL_3D_cuda.h b/src/integrators/VL_3D_cuda.h index 3310f5fe5..d3c58cc0b 100644 --- a/src/integrators/VL_3D_cuda.h +++ b/src/integrators/VL_3D_cuda.h @@ -10,8 +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); + Real dt, int n_fields, int custom_grav, Real density_floor, Real *host_grav_potential); void Free_Memory_VL_3D(); diff --git a/src/integrators/simple_3D_cuda.cu b/src/integrators/simple_3D_cuda.cu index 3a1ce77b1..8cbba2d07 100644 --- a/src/integrators/simple_3D_cuda.cu +++ b/src/integrators/simple_3D_cuda.cu @@ -26,8 +26,7 @@ void Simple_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) + Real dt, int n_fields, int custom_grav, Real density_floor, Real *host_grav_potential) { // Here, *dev_conserved contains the entire // set of conserved variables on the grid diff --git a/src/integrators/simple_3D_cuda.h b/src/integrators/simple_3D_cuda.h index 57d211ed1..67da69ca9 100644 --- a/src/integrators/simple_3D_cuda.h +++ b/src/integrators/simple_3D_cuda.h @@ -11,8 +11,7 @@ void Simple_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); + Real dt, int n_fields, int custom_grav, Real density_floor, Real *host_grav_potential); void Free_Memory_Simple_3D(); From cb5e8f3270926aecfec8f0d521afd4a27f552504 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Fri, 26 Jan 2024 10:21:02 -0500 Subject: [PATCH 30/40] run clang format --- 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 b32224c3e..8b72facdf 100644 --- a/src/dust/dust_cuda.cu +++ b/src/dust/dust_cuda.cu @@ -90,7 +90,7 @@ __global__ void Dust_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_g Real const temperature = hydro_utilities::Calc_Temp_Conserved(energy, density_gas, momentum_x, momentum_y, momentum_z, gamma, number_density); #endif // MHD - #endif // DE + #endif // DE Real tau_sp = Calc_Sputtering_Timescale(number_density, temperature, grain_radius) / TIME_UNIT; // sputtering timescale, kyr (sim units) From ddbd3125bc6227ee1a8f174f8f699fa94915b1a5 Mon Sep 17 00:00:00 2001 From: Bob Caddy Date: Thu, 18 Jan 2024 12:08:14 -0500 Subject: [PATCH 31/40] Update hydro build to use VL and PLMC Updated the default hydro build to use the Van Leer (VL) and PLMC build options. Also, updated the test data for those new builds --- builds/make.type.hydro | 6 +- cholla-tests-data | 2 +- src/reconstruction/plmc_cuda_tests.cu | 130 ++++++++++++++------------ 3 files changed, 74 insertions(+), 64 deletions(-) diff --git a/builds/make.type.hydro b/builds/make.type.hydro index f34d78172..b35dbd9ae 100644 --- a/builds/make.type.hydro +++ b/builds/make.type.hydro @@ -3,12 +3,12 @@ DFLAGS += -DCUDA DFLAGS += -DMPI_CHOLLA DFLAGS += -DPRECISION=2 -DFLAGS += -DPPMC +DFLAGS += -DPLMC DFLAGS += -DHLLC # Integrator -DFLAGS += -DSIMPLE -#DFLAGS += -DVL +# DFLAGS += -DSIMPLE +DFLAGS += -DVL # Apply a density and temperature floor DFLAGS += -DDENSITY_FLOOR diff --git a/cholla-tests-data b/cholla-tests-data index dcd73ff52..da5c3a309 160000 --- a/cholla-tests-data +++ b/cholla-tests-data @@ -1 +1 @@ -Subproject commit dcd73ff52b9027627b247c6d888bcdb56840c85e +Subproject commit da5c3a309d5451fabdec27fd7942e6121bb9c277 diff --git a/src/reconstruction/plmc_cuda_tests.cu b/src/reconstruction/plmc_cuda_tests.cu index 68f11b396..0207a09ac 100644 --- a/src/reconstruction/plmc_cuda_tests.cu +++ b/src/reconstruction/plmc_cuda_tests.cu @@ -25,6 +25,9 @@ TEST(tHYDROPlmcReconstructor, CorrectInputExpectCorrectOutput) { +#ifndef VL + #warning "The tHYDROPlmcReconstructor.CorrectInputExpectCorrectOutput only supports the Van Leer (VL) integrator" +#endif // VL // Set up PRNG to use std::mt19937_64 prng(42); std::uniform_real_distribution doubleRand(0.1, 5); @@ -49,66 +52,71 @@ TEST(tHYDROPlmcReconstructor, CorrectInputExpectCorrectOutput) dev_grid.cpyHostToDevice(host_grid); // Fiducial Data - std::vector> fiducial_interface_left = {{{26, 2.1584359129984056}, - {27, 0.70033864721549188}, - {106, 2.2476363309467553}, - {107, 3.0633780053857027}, - {186, 2.2245934101106259}, - {187, 2.1015872413794123}, - {266, 2.1263341057778309}, - {267, 3.9675148506537838}, - {346, 3.3640057502842691}, - {347, 21.091316282933843}}, - {{21, 0.72430827309279655}, - {37, 0.19457128219588618}, - {101, 5.4739527659741896}, - {117, 4.4286255636679313}, - {181, 0.12703829036056602}, - {197, 2.2851440769830953}, - {261, 1.5337035731959561}, - {277, 2.697375839048191}, - {341, 22.319601655044117}, - {357, 82.515887983144168}}, - {{25, 2.2863650183226212}, - {29, 1.686415421301841}, - {105, 0.72340346106443465}, - {109, 5.4713687086831388}, - {185, 3.929100145230096}, - {189, 4.9166140516911483}, - {265, 0.95177493689267167}, - {269, 0.46056494878491938}, - {345, 3.6886096301452787}, - {349, 16.105488797582133}}}; - std::vector> fiducial_interface_right = {{{25, 3.8877922383184833}, - {26, 0.70033864721549188}, - {105, 1.5947787943675635}, - {106, 3.0633780053857027}, - {185, 4.0069556576401011}, - {186, 2.1015872413794123}, - {265, 1.7883678016935785}, - {266, 3.9675148506537838}, - {345, 2.8032969746372527}, - {346, 21.091316282933843}}, - {{17, 0.43265217076853835}, - {33, 0.19457128219588618}, - {97, 3.2697645945288754}, - {113, 4.4286255636679313}, - {177, 0.07588397666718491}, - {193, 2.2851440769830953}, - {257, 0.91612950577699748}, - {273, 2.697375839048191}, - {337, 13.332201861384396}, - {353, 82.515887983144168}}, - {{5, 2.2863650183226212}, - {9, 1.686415421301841}, - {85, 0.72340346106443465}, - {89, 1.7792505446336098}, - {165, 5.3997753452111859}, - {169, 1.4379190463124139}, - {245, 0.95177493689267167}, - {249, 0.46056494878491938}, - {325, 6.6889498465051407}, - {329, 1.6145084086614281}}}; + std::vector> fiducial_interface_left = {{{26, 3.8877922383184833}, + {27, 0.70033864721549188}, + {106, 5.6625525038177784}, + {107, 3.0633780053857027}, + {186, 4.0069556576401011}, + {187, 2.1015872413794123}, + {266, 5.1729859852329314}, + {267, 3.9675148506537838}, + {346, 9.6301414677176531}, + {347, 21.091316282933843}}, + {{21, 0.74780807318015607}, + {37, 0.19457128219588618}, + {101, 5.6515522777659895}, + {117, 4.4286255636679313}, + {181, 0.13115998072061905}, + {197, 2.2851440769830953}, + {261, 1.5834637771067519}, + {277, 2.697375839048191}, + {341, 23.043749364531674}, + {357, 82.515887983144168}}, + {{25, 2.2863650183226212}, + {29, 1.686415421301841}, + {105, 0.72340346106443465}, + {109, 5.9563546443402542}, + {185, 3.6128571662018358}, + {189, 5.3735653401079038}, + {265, 0.95177493689267167}, + {269, 0.46056494878491938}, + {345, 3.1670194578067843}, + {349, 19.142817472509272}}}; + + std::vector> fiducial_interface_right = + + {{{25, 3.8877922383184833}, + {26, 0.70033864721549188}, + {105, 1.594778794367564}, + {106, 3.0633780053857027}, + {185, 4.0069556576401011}, + {186, 2.1015872413794123}, + {265, 1.7883678016935782}, + {266, 3.9675148506537838}, + {345, 2.8032969746372531}, + {346, 21.091316282933843}}, + {{17, 0.43265217076853835}, + {33, 0.19457128219588618}, + {97, 3.2697645945288754}, + {113, 4.4286255636679313}, + {177, 0.07588397666718491}, + {193, 2.2851440769830953}, + {257, 0.91612950577699748}, + {273, 2.697375839048191}, + {337, 13.332201861384396}, + {353, 82.515887983144168}}, + {{5, 2.2863650183226212}, + {9, 1.686415421301841}, + {85, 0.72340346106443465}, + {89, 1.77925054463361}, + {165, 5.3997753452111859}, + {169, 1.4379190463124141}, + {245, 0.95177493689267167}, + {249, 0.46056494878491938}, + {325, 6.6889498465051398}, + {329, 1.6145084086614285}}} + + ; // Loop over different directions for (size_t direction = 0; direction < 3; direction++) { @@ -161,6 +169,8 @@ TEST(tHYDROPlmcReconstructor, CorrectInputExpectCorrectOutput) ? 0.0 : fiducial_interface_right.at(direction)[i]; + // if (test_val != 0.0) std::cout << "{" << i << ", " << to_string_exact(test_val) << "}," << std::endl; + testing_utilities::Check_Results( fiducial_val, test_val, "right interface at i=" + std::to_string(i) + ", in direction " + std::to_string(direction)); From 29aa0519e798f9da107a477e694478ee7d1e842f Mon Sep 17 00:00:00 2001 From: Bob Caddy Date: Thu, 18 Jan 2024 13:36:52 -0500 Subject: [PATCH 32/40] Switch MHD to using PLMC by default Also updated test data --- builds/make.type.mhd | 2 +- src/system_tests/mhd_system_tests.cpp | 24 ++++++++++++------------ 2 files changed, 13 insertions(+), 13 deletions(-) diff --git a/builds/make.type.mhd b/builds/make.type.mhd index 4459819e8..d08e6373e 100644 --- a/builds/make.type.mhd +++ b/builds/make.type.mhd @@ -9,7 +9,7 @@ MPI_GPU ?= DFLAGS += -DCUDA DFLAGS += -DMPI_CHOLLA DFLAGS += -DPRECISION=2 -DFLAGS += -DPPMC +DFLAGS += -DPLMC DFLAGS += -DHLLD DFLAGS += -DMHD diff --git a/src/system_tests/mhd_system_tests.cpp b/src/system_tests/mhd_system_tests.cpp index 4261797b2..a14caa9a1 100644 --- a/src/system_tests/mhd_system_tests.cpp +++ b/src/system_tests/mhd_system_tests.cpp @@ -228,7 +228,7 @@ TEST_P(tMHDSYSTEMLinearWavesParameterizedAngle, SlowMagnetosonicWaveRightMovingC #ifdef PCM waveTest.runL1ErrorTest(4.E-7, 4.E-7); #elif defined(PLMC) - waveTest.runL1ErrorTest(2.0E-8, 2.7E-8); + waveTest.runL1ErrorTest(2.0E-8, 2.75E-8); #elif defined(PPMC) waveTest.runL1ErrorTest(1.45E-9, 1.3E-9); #endif // PCM @@ -266,7 +266,7 @@ TEST_P(tMHDSYSTEMLinearWavesParameterizedAngle, SlowMagnetosonicWaveLeftMovingCo #ifdef PCM waveTest.runL1ErrorTest(4.E-7, 4.E-7); #elif defined(PLMC) - waveTest.runL1ErrorTest(2.0E-8, 2.7E-8); + waveTest.runL1ErrorTest(2.0E-8, 2.75E-8); #elif defined(PPMC) waveTest.runL1ErrorTest(1.45E-9, 1.3E-9); #endif // PCM @@ -416,12 +416,12 @@ TEST_P(tMHDSYSTEMLinearWavesParameterizedAngle, FastMagnetosonicWaveExpectSecond waveTest.setFiducialNumTimeSteps(numTimeSteps[domain_direction - 1]); // Run the wave - waveTest.runL1ErrorTest(3.0E-8, 4.0E-8); + waveTest.runL1ErrorTest(7.0E-8, 1.5E-7); // Check the scaling double const low_res_l2norm = waveTest.getL2Norm(); testing_utilities::Check_Results(4.0, low_res_l2norm / high_res_l2norms["fast_" + std::to_string(domain_direction)], - "", 0.17); + "", 0.2); } TEST_P(tMHDSYSTEMLinearWavesParameterizedAngle, SlowMagnetosonicWaveExpectSecondOrderConvergence) @@ -452,12 +452,12 @@ TEST_P(tMHDSYSTEMLinearWavesParameterizedAngle, SlowMagnetosonicWaveExpectSecond waveTest.setFiducialNumTimeSteps(numTimeSteps[domain_direction - 1]); // Run the wave - waveTest.runL1ErrorTest(3.0E-8, 4.0E-8); + waveTest.runL1ErrorTest(5.4E-8, 8.0E-8); // Check the scaling double const low_res_l2norm = waveTest.getL2Norm(); testing_utilities::Check_Results(4.0, low_res_l2norm / high_res_l2norms["slow_" + std::to_string(domain_direction)], - "", 0.17); + "", 0.2); } TEST_P(tMHDSYSTEMLinearWavesParameterizedAngle, AlfvenWaveExpectSecondOrderConvergence) @@ -487,12 +487,12 @@ TEST_P(tMHDSYSTEMLinearWavesParameterizedAngle, AlfvenWaveExpectSecondOrderConve waveTest.setFiducialNumTimeSteps(numTimeSteps[domain_direction - 1]); // Run the wave - waveTest.runL1ErrorTest(3.0E-8, 4.0E-8); + waveTest.runL1ErrorTest(4.5E-8, 8.0E-8); // Check the scaling double const low_res_l2norm = waveTest.getL2Norm(); testing_utilities::Check_Results(4.0, low_res_l2norm / high_res_l2norms["alfven_" + std::to_string(domain_direction)], - "", 0.17); + "", 0.2); } TEST_P(tMHDSYSTEMLinearWavesParameterizedAngle, MHDContactWaveExpectSecondOrderConvergence) @@ -523,12 +523,12 @@ TEST_P(tMHDSYSTEMLinearWavesParameterizedAngle, MHDContactWaveExpectSecondOrderC waveTest.setFiducialNumTimeSteps(numTimeSteps[domain_direction - 1]); // Run the wave - waveTest.runL1ErrorTest(3.0E-8, 4.0E-8); + waveTest.runL1ErrorTest(5.0E-8, 8.0E-8); // Check the scaling double const low_res_l2norm = waveTest.getL2Norm(); testing_utilities::Check_Results( - 4.0, low_res_l2norm / high_res_l2norms["contact_" + std::to_string(domain_direction)], "", 0.17); + 4.0, low_res_l2norm / high_res_l2norms["contact_" + std::to_string(domain_direction)], "", 0.2); } INSTANTIATE_TEST_SUITE_P(, tMHDSYSTEMLinearWavesParameterizedAngle, @@ -645,7 +645,7 @@ TEST_P(tMHDSYSTEMLinearWavesParameterizedMpi, SlowMagnetosonicWaveRightMovingCor #ifdef PCM waveTest.runL1ErrorTest(4.E-7, 4.E-7); #elif defined(PLMC) - waveTest.runL1ErrorTest(2.0E-8, 2.7E-8); + waveTest.runL1ErrorTest(2.0E-8, 2.75E-8); #elif defined(PPMC) waveTest.runL1ErrorTest(1.4E-9, 1.3E-9); #endif // PCM @@ -681,7 +681,7 @@ TEST_P(tMHDSYSTEMLinearWavesParameterizedMpi, SlowMagnetosonicWaveLeftMovingCorr #ifdef PCM waveTest.runL1ErrorTest(4.E-7, 4.E-7); #elif defined(PLMC) - waveTest.runL1ErrorTest(2.0E-8, 2.7E-8); + waveTest.runL1ErrorTest(2.0E-8, 2.8E-8); #elif defined(PPMC) waveTest.runL1ErrorTest(1.4E-9, 1.3E-9); #endif // PCM From fafcf4f3946fd8856ce8bad6ac44b1d31faabc3c Mon Sep 17 00:00:00 2001 From: Bob Caddy Date: Thu, 18 Jan 2024 13:40:23 -0500 Subject: [PATCH 33/40] Update other builds to PLMC and VL default Updated template, basic_scalar, disk, dust, rot_proj, static_grav to use PLMC and Van Leer integrator --- builds/make.inc.template | 4 ++-- builds/make.type.basic_scalar | 6 +++--- builds/make.type.disk | 2 +- builds/make.type.dust | 2 +- builds/make.type.rot_proj | 6 +++--- builds/make.type.static_grav | 6 +++--- 6 files changed, 13 insertions(+), 13 deletions(-) diff --git a/builds/make.inc.template b/builds/make.inc.template index 3ae156225..98d2c146c 100644 --- a/builds/make.inc.template +++ b/builds/make.inc.template @@ -33,8 +33,8 @@ DFLAGS += -DHDF5 # Reconstruction #DFLAGS += -DPCM #DFLAGS += -DPLMP -#DFLAGS += -DPLMC -DFLAGS += -DPPMP +DFLAGS += -DPLMC +#DFLAGS += -DPPMP #DFLAGS += -DPPMC # Riemann Solver diff --git a/builds/make.type.basic_scalar b/builds/make.type.basic_scalar index d2dd75892..5aa4a5d0e 100644 --- a/builds/make.type.basic_scalar +++ b/builds/make.type.basic_scalar @@ -3,12 +3,12 @@ DFLAGS += -DCUDA DFLAGS += -DMPI_CHOLLA DFLAGS += -DPRECISION=2 -DFLAGS += -DPPMC +DFLAGS += -DPLMC DFLAGS += -DHLLC # Integrator -DFLAGS += -DSIMPLE -#DFLAGS += -DVL +# DFLAGS += -DSIMPLE +DFLAGS += -DVL # Apply a density and temperature floor DFLAGS += -DDENSITY_FLOOR diff --git a/builds/make.type.disk b/builds/make.type.disk index f2e3f0ec1..284b2c73d 100644 --- a/builds/make.type.disk +++ b/builds/make.type.disk @@ -24,7 +24,7 @@ DFLAGS += -DGRAVITY_5_POINTS_GRADIENT DFLAGS += -DCUDA DFLAGS += -DMPI_CHOLLA DFLAGS += -DPRECISION=2 -DFLAGS += -DPPMC +DFLAGS += -DPLMC DFLAGS += -DHLLC DFLAGS += -DVL diff --git a/builds/make.type.dust b/builds/make.type.dust index 24a765302..5addcacdc 100644 --- a/builds/make.type.dust +++ b/builds/make.type.dust @@ -9,7 +9,7 @@ MPI_GPU ?= DFLAGS += -DCUDA DFLAGS += -DMPI_CHOLLA DFLAGS += -DPRECISION=2 -DFLAGS += -DPLMP +DFLAGS += -DPLMC DFLAGS += -DHLLC DFLAGS += -DDE diff --git a/builds/make.type.rot_proj b/builds/make.type.rot_proj index e6faa7514..e29ab43e7 100644 --- a/builds/make.type.rot_proj +++ b/builds/make.type.rot_proj @@ -3,12 +3,12 @@ DFLAGS += -DCUDA DFLAGS += -DMPI_CHOLLA DFLAGS += -DPRECISION=2 -DFLAGS += -DPPMC +DFLAGS += -DPLMC DFLAGS += -DHLLC # Integrator -DFLAGS += -DSIMPLE -#DFLAGS += -DVL +# DFLAGS += -DSIMPLE +DFLAGS += -DVL # Apply a density and temperature floor DFLAGS += -DDENSITY_FLOOR diff --git a/builds/make.type.static_grav b/builds/make.type.static_grav index 4f13e7288..cd77643f2 100644 --- a/builds/make.type.static_grav +++ b/builds/make.type.static_grav @@ -3,12 +3,12 @@ DFLAGS += -DCUDA DFLAGS += -DMPI_CHOLLA DFLAGS += -DPRECISION=2 -DFLAGS += -DPPMC +DFLAGS += -DPLMC DFLAGS += -DHLLC # Integrator -DFLAGS += -DSIMPLE -#DFLAGS += -DVL +# DFLAGS += -DSIMPLE +DFLAGS += -DVL # Apply a density and temperature floor DFLAGS += -DDENSITY_FLOOR From 2215b2f643286f367458ce2c195ffab09a0b1a1b Mon Sep 17 00:00:00 2001 From: Bob Caddy Date: Mon, 22 Jan 2024 13:54:24 -0500 Subject: [PATCH 34/40] Make sure cosmology builds use SIMPLE integrator --- builds/make.type.hydro | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/builds/make.type.hydro b/builds/make.type.hydro index b35dbd9ae..1a96baaa9 100644 --- a/builds/make.type.hydro +++ b/builds/make.type.hydro @@ -7,8 +7,11 @@ DFLAGS += -DPLMC DFLAGS += -DHLLC # Integrator -# DFLAGS += -DSIMPLE +ifeq ($(findstring cosmology,$(TYPE)),cosmology) +DFLAGS += -DSIMPLE +else DFLAGS += -DVL +endif # Apply a density and temperature floor DFLAGS += -DDENSITY_FLOOR From ff09d5dc0939b98710652066a97610e6a5472c3d Mon Sep 17 00:00:00 2001 From: Bob Caddy Date: Tue, 23 Jan 2024 17:02:12 -0500 Subject: [PATCH 35/40] Fix clang-tidy warning --- src/global/global.cpp | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/src/global/global.cpp b/src/global/global.cpp index 1a867965c..2d3a23467 100644 --- a/src/global/global.cpp +++ b/src/global/global.cpp @@ -242,6 +242,10 @@ void Parse_Param(char *name, char *value, struct Parameters *parms) #endif // DE } else if (strcmp(name, "output_always") == 0) { int tmp = atoi(value); + // In this case the CHOLLA_ASSERT macro runs into issuse with the readability-simplify-boolean-expr clang-tidy check + // due to some weird macro expansion stuff. That check has been disabled here for now but in clang-tidy 18 the + // IgnoreMacro option should be used instead. + // NOLINTNEXTLINE(readability-simplify-boolean-expr) CHOLLA_ASSERT((tmp == 0) or (tmp == 1), "output_always must be 1 or 0."); parms->output_always = tmp; #ifdef MHD From b4589620da6750efceff188d213c6e46cdc4a8c5 Mon Sep 17 00:00:00 2001 From: Bob Caddy Date: Tue, 23 Jan 2024 17:29:13 -0500 Subject: [PATCH 36/40] Replace #warning with print to cerr Avoids issues with running clang-tidy in non-VL builds --- src/reconstruction/plmc_cuda_tests.cu | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/src/reconstruction/plmc_cuda_tests.cu b/src/reconstruction/plmc_cuda_tests.cu index 0207a09ac..678f6329d 100644 --- a/src/reconstruction/plmc_cuda_tests.cu +++ b/src/reconstruction/plmc_cuda_tests.cu @@ -26,7 +26,10 @@ TEST(tHYDROPlmcReconstructor, CorrectInputExpectCorrectOutput) { #ifndef VL - #warning "The tHYDROPlmcReconstructor.CorrectInputExpectCorrectOutput only supports the Van Leer (VL) integrator" + std::cerr << "Warning: The tHYDROPlmcReconstructor.CorrectInputExpectCorrectOutput only supports the Van Leer (VL) " + "integrator" + << std::endl; + return; #endif // VL // Set up PRNG to use std::mt19937_64 prng(42); From 3e1c1959f277272cb17c362ca1c4169d36ef8567 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Thu, 1 Feb 2024 14:05:14 -0500 Subject: [PATCH 37/40] set default floor values in global.h and add warning --- src/global/global.cpp | 9 +++++++++ src/global/global.h | 6 +++--- 2 files changed, 12 insertions(+), 3 deletions(-) diff --git a/src/global/global.cpp b/src/global/global.cpp index 2d3a23467..49bc5f681 100644 --- a/src/global/global.cpp +++ b/src/global/global.cpp @@ -439,14 +439,23 @@ void Parse_Param(char *name, char *value, struct Parameters *parms) #ifdef TEMPERATURE_FLOOR } else if (strcmp(name, "temperature_floor") == 0) { parms->temperature_floor = atof(value); + if (parms->temperature_floor == 0) { + chprintf("WARNING: temperature floor is set to its default value!\n"); + } #endif #ifdef DENSITY_FLOOR } else if (strcmp(name, "density_floor") == 0) { parms->density_floor = atof(value); + if (parms->density_floor == 0) { + chprintf("WARNING: density floor is set to its default value!\n"); + } #endif #ifdef SCALAR_FLOOR } else if (strcmp(name, "scalar_floor") == 0) { parms->scalar_floor = atof(value); + if (parms->scalar_floor == 0) { + chprintf("WARNING: scalar floor is set to its default value!\n"); + } #endif #ifdef ANALYSIS } else if (strcmp(name, "analysis_scale_outputs_file") == 0) { diff --git a/src/global/global.h b/src/global/global.h index 3cf58312a..8fd8b505b 100644 --- a/src/global/global.h +++ b/src/global/global.h @@ -323,13 +323,13 @@ struct Parameters { // photoionization rates of HI, HeI and HeII #endif #ifdef TEMPERATURE_FLOOR - Real temperature_floor; + Real temperature_floor = 0; #endif #ifdef DENSITY_FLOOR - Real density_floor; + Real density_floor = 0; #endif #ifdef SCALAR_FLOOR - Real scalar_floor; + Real scalar_floor = 0; #endif #ifdef ANALYSIS char analysis_scale_outputs_file[MAXLEN]; // File for the scale_factor output From 114f96b555d3ac01822ba6925bdb124d6c17487b Mon Sep 17 00:00:00 2001 From: helenarichie Date: Thu, 1 Feb 2024 14:56:15 -0500 Subject: [PATCH 38/40] change warning message and tweak how the default floors are set --- src/global/global.cpp | 6 +++--- src/global/global.h | 6 ------ src/grid/grid3D.cpp | 6 ------ 3 files changed, 3 insertions(+), 15 deletions(-) diff --git a/src/global/global.cpp b/src/global/global.cpp index 49bc5f681..ed140059c 100644 --- a/src/global/global.cpp +++ b/src/global/global.cpp @@ -440,21 +440,21 @@ void Parse_Param(char *name, char *value, struct Parameters *parms) } else if (strcmp(name, "temperature_floor") == 0) { parms->temperature_floor = atof(value); if (parms->temperature_floor == 0) { - chprintf("WARNING: temperature floor is set to its default value!\n"); + chprintf("WARNING: temperature floor is set to its default value (zero)! It can be set to a different value in the input parameter file.\n"); } #endif #ifdef DENSITY_FLOOR } else if (strcmp(name, "density_floor") == 0) { parms->density_floor = atof(value); if (parms->density_floor == 0) { - chprintf("WARNING: density floor is set to its default value!\n"); + chprintf("WARNING: density floor is set to its default value (zero)! It can be set to a different value in the input parameter file.\n"); } #endif #ifdef SCALAR_FLOOR } else if (strcmp(name, "scalar_floor") == 0) { parms->scalar_floor = atof(value); if (parms->scalar_floor == 0) { - chprintf("WARNING: scalar floor is set to its default value!\n"); + chprintf("WARNING: scalar floor is set to its default value (zero)! It can be set to a different value in the input parameter file.\n"); } #endif #ifdef ANALYSIS diff --git a/src/global/global.h b/src/global/global.h index 8fd8b505b..91c54911e 100644 --- a/src/global/global.h +++ b/src/global/global.h @@ -322,15 +322,9 @@ struct Parameters { char UVB_rates_file[MAXLEN]; // File for the UVB photoheating and // photoionization rates of HI, HeI and HeII #endif -#ifdef TEMPERATURE_FLOOR Real temperature_floor = 0; -#endif -#ifdef DENSITY_FLOOR Real density_floor = 0; -#endif -#ifdef SCALAR_FLOOR Real scalar_floor = 0; -#endif #ifdef ANALYSIS char analysis_scale_outputs_file[MAXLEN]; // File for the scale_factor output // values for cosmological diff --git a/src/grid/grid3D.cpp b/src/grid/grid3D.cpp index 1cfde07a6..b773e18ea 100644 --- a/src/grid/grid3D.cpp +++ b/src/grid/grid3D.cpp @@ -261,20 +261,14 @@ void Grid3D::Initialize(struct Parameters *P) #ifdef TEMPERATURE_FLOOR H.temperature_floor = P->temperature_floor; #else - H.temperature_floor = 0.0; -#endif #ifdef DENSITY_FLOOR H.density_floor = P->density_floor; #else - H.density_floor = 0.0; -#endif #ifdef SCALAR_FLOOR H.scalar_floor = P->scalar_floor; #else - H.scalar_floor = 0.0; -#endif #ifdef COSMOLOGY H.OUTPUT_SCALE_FACOR = not(P->scale_outputs_file[0] == '\0'); From 3eb2fb8eb19825ed653792cf552dfa66215cf9f4 Mon Sep 17 00:00:00 2001 From: helenarichie Date: Thu, 1 Feb 2024 15:16:59 -0500 Subject: [PATCH 39/40] fix syntax errors --- src/grid/grid3D.cpp | 8 +++----- src/hydro/hydro_cuda.cu | 2 -- 2 files changed, 3 insertions(+), 7 deletions(-) diff --git a/src/grid/grid3D.cpp b/src/grid/grid3D.cpp index b773e18ea..9ca5ba69a 100644 --- a/src/grid/grid3D.cpp +++ b/src/grid/grid3D.cpp @@ -260,15 +260,15 @@ void Grid3D::Initialize(struct Parameters *P) // Values for lower limit for density and temperature #ifdef TEMPERATURE_FLOOR H.temperature_floor = P->temperature_floor; -#else +#endif #ifdef DENSITY_FLOOR H.density_floor = P->density_floor; -#else +#endif #ifdef SCALAR_FLOOR H.scalar_floor = P->scalar_floor; -#else +#endif #ifdef COSMOLOGY H.OUTPUT_SCALE_FACOR = not(P->scale_outputs_file[0] == '\0'); @@ -490,8 +490,6 @@ Real Grid3D::Update_Hydro_Grid() Execute_Hydro_Integrator(); -#ifdef CUDA - #ifdef TEMPERATURE_FLOOR // Set the lower limit temperature (Internal Energy) Real U_floor; diff --git a/src/hydro/hydro_cuda.cu b/src/hydro/hydro_cuda.cu index 8a69266ff..4c45c87f2 100644 --- a/src/hydro/hydro_cuda.cu +++ b/src/hydro/hydro_cuda.cu @@ -1305,5 +1305,3 @@ __global__ void Scalar_Floor_Kernel(Real *dev_conserved, int nx, int ny, int nz, } } } - -#endif // CUDA From cede7f3a1b65b149a74f35a7da96fe24478f76ac Mon Sep 17 00:00:00 2001 From: helenarichie Date: Thu, 1 Feb 2024 15:19:06 -0500 Subject: [PATCH 40/40] run clang format --- src/global/global.cpp | 12 +++++++++--- src/global/global.h | 4 ++-- src/grid/grid3D.cpp | 24 ++++++++++++------------ src/hydro/hydro_cuda.cu | 4 ++-- 4 files changed, 25 insertions(+), 19 deletions(-) diff --git a/src/global/global.cpp b/src/global/global.cpp index ed140059c..a5e4068a4 100644 --- a/src/global/global.cpp +++ b/src/global/global.cpp @@ -440,21 +440,27 @@ void Parse_Param(char *name, char *value, struct Parameters *parms) } else if (strcmp(name, "temperature_floor") == 0) { parms->temperature_floor = atof(value); if (parms->temperature_floor == 0) { - chprintf("WARNING: temperature floor is set to its default value (zero)! It can be set to a different value in the input parameter file.\n"); + chprintf( + "WARNING: temperature floor is set to its default value (zero)! It can be set to a different value in the " + "input parameter file.\n"); } #endif #ifdef DENSITY_FLOOR } else if (strcmp(name, "density_floor") == 0) { parms->density_floor = atof(value); if (parms->density_floor == 0) { - chprintf("WARNING: density floor is set to its default value (zero)! It can be set to a different value in the input parameter file.\n"); + chprintf( + "WARNING: density floor is set to its default value (zero)! It can be set to a different value in the input " + "parameter file.\n"); } #endif #ifdef SCALAR_FLOOR } else if (strcmp(name, "scalar_floor") == 0) { parms->scalar_floor = atof(value); if (parms->scalar_floor == 0) { - chprintf("WARNING: scalar floor is set to its default value (zero)! It can be set to a different value in the input parameter file.\n"); + chprintf( + "WARNING: scalar floor is set to its default value (zero)! It can be set to a different value in the input " + "parameter file.\n"); } #endif #ifdef ANALYSIS diff --git a/src/global/global.h b/src/global/global.h index 91c54911e..b1a906532 100644 --- a/src/global/global.h +++ b/src/global/global.h @@ -323,8 +323,8 @@ struct Parameters { // photoionization rates of HI, HeI and HeII #endif Real temperature_floor = 0; - Real density_floor = 0; - Real scalar_floor = 0; + Real density_floor = 0; + Real scalar_floor = 0; #ifdef ANALYSIS char analysis_scale_outputs_file[MAXLEN]; // File for the scale_factor output // values for cosmological diff --git a/src/grid/grid3D.cpp b/src/grid/grid3D.cpp index 9ca5ba69a..ef4d57928 100644 --- a/src/grid/grid3D.cpp +++ b/src/grid/grid3D.cpp @@ -458,7 +458,7 @@ void Grid3D::Execute_Hydro_Integrator(void) 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, H.dz, H.xbound, H.ybound, H.zbound, H.dt, H.n_fields, H.custom_grav, H.density_floor, C.Grav_potential); - #endif // SIMPLE +#endif // SIMPLE } else { chprintf("Error: Grid dimensions nx: %d ny: %d nz: %d not supported.\n", H.nx, H.ny, H.nz); chexit(-1); @@ -490,27 +490,27 @@ Real Grid3D::Update_Hydro_Grid() Execute_Hydro_Integrator(); - #ifdef TEMPERATURE_FLOOR +#ifdef TEMPERATURE_FLOOR // Set the lower limit temperature (Internal Energy) Real U_floor; // Minimum of internal energy from minumum of temperature U_floor = H.temperature_floor * KB / (gama - 1) / MP / SP_ENERGY_UNIT; - #ifdef COSMOLOGY + #ifdef COSMOLOGY U_floor = H.temperature_floor / (gama - 1) / MP * KB * 1e-10; // ( km/s )^2 U_floor /= Cosmo.v_0_gas * Cosmo.v_0_gas / Cosmo.current_a / Cosmo.current_a; - #endif + #endif Apply_Temperature_Floor(C.device, H.nx, H.ny, H.nz, H.n_ghost, H.n_fields, U_floor); - #endif // TEMPERATURE_FLOOR +#endif // TEMPERATURE_FLOOR - #ifdef SCALAR_FLOOR - #ifdef DUST +#ifdef SCALAR_FLOOR + #ifdef DUST Apply_Scalar_Floor(C.device, H.nx, H.ny, H.nz, H.n_ghost, grid_enum::dust_density, H.scalar_floor); - #endif - #endif // SCALAR_FLOOR + #endif +#endif // SCALAR_FLOOR - // == Perform chemistry/cooling (there are a few different cases) == - #ifdef COOLING_GPU - #ifdef CPU_TIME +// == Perform chemistry/cooling (there are a few different cases) == +#ifdef COOLING_GPU + #ifdef CPU_TIME Timer.Cooling_GPU.Start(); #endif // ==Apply Cooling from cooling/cooling_cuda.h== diff --git a/src/hydro/hydro_cuda.cu b/src/hydro/hydro_cuda.cu index 4c45c87f2..4d0661fbd 100644 --- a/src/hydro/hydro_cuda.cu +++ b/src/hydro/hydro_cuda.cu @@ -1143,12 +1143,12 @@ __global__ void Temperature_Floor_Kernel(Real *dev_conserved, int nx, int ny, in dev_conserved[4 * n_cells + id] = Ekin + d * U_floor; } - #ifdef DE +#ifdef DE U = dev_conserved[(n_fields - 1) * n_cells + id] / d; if (U < U_floor) { dev_conserved[(n_fields - 1) * n_cells + id] = d * U_floor; } - #endif +#endif } }