From 39a6f977feedffe98eacf8dc11cb9493cf8150a9 Mon Sep 17 00:00:00 2001 From: Evan Schneider Date: Fri, 8 Dec 2023 14:53:05 -0500 Subject: [PATCH 1/6] Fixed the way cell averaging works so that conserved values are internally consistent when averaging all fields. --- src/hydro/hydro_cuda.cu | 100 ++++++++++++++++++++++++++++++++++------ src/hydro/hydro_cuda.h | 2 +- 2 files changed, 87 insertions(+), 15 deletions(-) diff --git a/src/hydro/hydro_cuda.cu b/src/hydro/hydro_cuda.cu index c74654c0e..983a6e506 100644 --- a/src/hydro/hydro_cuda.cu +++ b/src/hydro/hydro_cuda.cu @@ -1153,23 +1153,95 @@ __device__ Real Average_Cell_Single_Field(int field_indx, int i, int j, int k, i } __device__ void Average_Cell_All_Fields(int i, int j, int k, int nx, int ny, int nz, int ncells, int n_fields, - Real *conserved) + Real gamma, Real *conserved) { - // Average Density - Average_Cell_Single_Field(0, i, j, k, nx, ny, nz, ncells, conserved); - // Average Momentum_x - Average_Cell_Single_Field(1, i, j, k, nx, ny, nz, ncells, conserved); - // Average Momentum_y - Average_Cell_Single_Field(2, i, j, k, nx, ny, nz, ncells, conserved); - // Average Momentum_z - Average_Cell_Single_Field(3, i, j, k, nx, ny, nz, ncells, conserved); - // Average Energy - Average_Cell_Single_Field(4, i, j, k, nx, ny, nz, ncells, conserved); + int id = i + (j)*nx + (k)*nx*ny; + + Real d, mx, my, mz, E, P; + d = conserved[grid_enum::density*ncells + id]; + mx = conserved[grid_enum::momentumx*ncells + id]; + my = conserved[grid_enum::momentumy*ncells + id]; + mz = conserved[grid_enum::momentumz*ncells + id]; + E = conserved[grid_enum::Energy*ncells + id]; + P = (E - (0.5/d)*(mx*mx + my*my + mz*mz))*(gamma-1.0); + + printf("%3d %3d %3d BC: d: %e E:%e P:%e vx:%e vy:%e vz:%e\n", i, j, k, d, E, P, mx/d, my/d, mz/d); + + int idn; + int N = 0; + Real d_av, vx_av, vy_av, vz_av, P_av; + d_av = vx_av = vy_av = vz_av = P_av = 0.0; + #ifdef SCALAR + Real scalar[NSCALARS], scalar_av[NSCALARS]; + for (int n=0; n 0.0 && P > 0.0) { + d_av += d; + vx_av += mx; + vy_av += my; + vz_av += mz; + P_av += P/(gamma-1.0); + #ifdef SCALAR + for (int n=0; n Date: Fri, 8 Dec 2023 15:45:24 -0500 Subject: [PATCH 2/6] clang formatted and compiling --- src/hydro/hydro_cuda.cu | 93 ++++++++++++++++++++--------------------- 1 file changed, 46 insertions(+), 47 deletions(-) diff --git a/src/hydro/hydro_cuda.cu b/src/hydro/hydro_cuda.cu index 983a6e506..40f89a347 100644 --- a/src/hydro/hydro_cuda.cu +++ b/src/hydro/hydro_cuda.cu @@ -1155,17 +1155,17 @@ __device__ Real Average_Cell_Single_Field(int field_indx, int i, int j, int k, i __device__ void Average_Cell_All_Fields(int i, int j, int k, int nx, int ny, int nz, int ncells, int n_fields, Real gamma, Real *conserved) { - int id = i + (j)*nx + (k)*nx*ny; + int id = i + (j)*nx + (k)*nx * ny; Real d, mx, my, mz, E, P; - d = conserved[grid_enum::density*ncells + id]; - mx = conserved[grid_enum::momentumx*ncells + id]; - my = conserved[grid_enum::momentumy*ncells + id]; - mz = conserved[grid_enum::momentumz*ncells + id]; - E = conserved[grid_enum::Energy*ncells + id]; - P = (E - (0.5/d)*(mx*mx + my*my + mz*mz))*(gamma-1.0); + d = conserved[grid_enum::density * ncells + id]; + mx = conserved[grid_enum::momentum_x * ncells + id]; + my = conserved[grid_enum::momentum_y * ncells + id]; + mz = conserved[grid_enum::momentum_z * ncells + id]; + E = conserved[grid_enum::Energy * ncells + id]; + P = (E - (0.5 / d) * (mx * mx + my * my + mz * mz)) * (gamma - 1.0); - printf("%3d %3d %3d BC: d: %e E:%e P:%e vx:%e vy:%e vz:%e\n", i, j, k, d, E, P, mx/d, my/d, mz/d); + printf("%3d %3d %3d BC: d: %e E:%e P:%e vx:%e vy:%e vz:%e\n", i, j, k, d, E, P, mx / d, my / d, mz / d); int idn; int N = 0; @@ -1173,75 +1173,74 @@ __device__ void Average_Cell_All_Fields(int i, int j, int k, int nx, int ny, int d_av = vx_av = vy_av = vz_av = P_av = 0.0; #ifdef SCALAR Real scalar[NSCALARS], scalar_av[NSCALARS]; - for (int n=0; n 0.0 && P > 0.0) { d_av += d; vx_av += mx; vy_av += my; vz_av += mz; - P_av += P/(gamma-1.0); - #ifdef SCALAR - for (int n=0; n Date: Fri, 8 Dec 2023 15:50:23 -0500 Subject: [PATCH 3/6] added average cells to cell update crash --- src/hydro/hydro_cuda.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/hydro/hydro_cuda.cu b/src/hydro/hydro_cuda.cu index 40f89a347..854c11147 100644 --- a/src/hydro/hydro_cuda.cu +++ b/src/hydro/hydro_cuda.cu @@ -385,6 +385,7 @@ __global__ void Update_Conserved_Variables_3D(Real *dev_conserved, Real *Q_Lx, R printf("%3d %3d %3d Thread crashed in final update. %e %e %e %e %e\n", xid + x_off, yid + y_off, zid + z_off, dev_conserved[id], dtodx * (dev_F_x[imo] - dev_F_x[id]), dtody * (dev_F_y[jmo] - dev_F_y[id]), dtodz * (dev_F_z[kmo] - dev_F_z[id]), dev_conserved[4 * n_cells + id]); + Average_Cell_All_Fields(xid, yid, zid, nx, ny, nz, n_cells, n_fields, gamma, dev_conserved); } #endif // DENSITY_FLOOR /* @@ -653,7 +654,7 @@ __global__ void Average_Slow_Cells_3D(Real *dev_conserved, int nx, int ny, int n xid, yid, zid, 1. / max_dti, 1. / max_dti_slow, dev_conserved[id] * DENSITY_UNIT / 0.6 / MP, temp, speed * VELOCITY_UNIT * 1e-5, vx * VELOCITY_UNIT * 1e-5, vy * VELOCITY_UNIT * 1e-5, vz * VELOCITY_UNIT * 1e-5, cs); - Average_Cell_All_Fields(xid, yid, zid, nx, ny, nz, n_cells, n_fields, dev_conserved); + Average_Cell_All_Fields(xid, yid, zid, nx, ny, nz, n_cells, n_fields, gamma, dev_conserved); } } } From f5736441a743eb2d0d140e17a0fadc80584973cb Mon Sep 17 00:00:00 2001 From: Evan Schneider Date: Mon, 11 Dec 2023 13:20:56 -0500 Subject: [PATCH 4/6] turned off linting for scalar for loop --- src/hydro/hydro_cuda.cu | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/src/hydro/hydro_cuda.cu b/src/hydro/hydro_cuda.cu index d2b7a666d..ac914cfdf 100644 --- a/src/hydro/hydro_cuda.cu +++ b/src/hydro/hydro_cuda.cu @@ -1174,7 +1174,7 @@ __device__ void Average_Cell_All_Fields(int i, int j, int k, int nx, int ny, int d_av = vx_av = vy_av = vz_av = P_av = 0.0; #ifdef SCALAR Real scalar[NSCALARS], scalar_av[NSCALARS]; - for (int n = 0; n < NSCALARS; n++) { + for (int n = 0; n < NSCALARS; n++) { //NOLINT scalar_av[n] = 0.0; } #endif @@ -1189,7 +1189,7 @@ __device__ void Average_Cell_All_Fields(int i, int j, int k, int nx, int ny, int mz = conserved[grid_enum::momentum_z * ncells + idn]; P = (conserved[grid_enum::Energy * ncells + idn] - (0.5 / d) * (mx * mx + my * my + mz * mz)) * (gamma - 1.0); #ifdef SCALAR - for (int n = 0; n < NSCALARS; n++) { + for (int n = 0; n < NSCALARS; n++) { //NOLINT scalar[n] = conserved[grid_enum::scalar * ncells + idn]; } #endif @@ -1200,7 +1200,7 @@ __device__ void Average_Cell_All_Fields(int i, int j, int k, int nx, int ny, int vz_av += mz; P_av += P / (gamma - 1.0); #ifdef SCALAR - for (int n = 0; n < NSCALARS; n++) { + for (int n = 0; n < NSCALARS; n++) { //NOLINT scalar_av[n] += scalar[n]; } #endif @@ -1215,7 +1215,7 @@ __device__ void Average_Cell_All_Fields(int i, int j, int k, int nx, int ny, int vy_av = vy_av / d_av; vz_av = vz_av / d_av; #ifdef SCALAR - for (int n = 0; n < NSCALARS; n++) { + for (int n = 0; n < NSCALARS; n++) { //NOLINT scalar_av[n] = scalar_av[n] / d_av; } #endif From bcd3a2770f12be1f0333a5dc6449042f640c2fe3 Mon Sep 17 00:00:00 2001 From: Evan Schneider Date: Mon, 11 Dec 2023 13:24:46 -0500 Subject: [PATCH 5/6] clang format --- src/hydro/hydro_cuda.cu | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/src/hydro/hydro_cuda.cu b/src/hydro/hydro_cuda.cu index ac914cfdf..0127344af 100644 --- a/src/hydro/hydro_cuda.cu +++ b/src/hydro/hydro_cuda.cu @@ -1174,7 +1174,7 @@ __device__ void Average_Cell_All_Fields(int i, int j, int k, int nx, int ny, int d_av = vx_av = vy_av = vz_av = P_av = 0.0; #ifdef SCALAR Real scalar[NSCALARS], scalar_av[NSCALARS]; - for (int n = 0; n < NSCALARS; n++) { //NOLINT + for (int n = 0; n < NSCALARS; n++) { // NOLINT scalar_av[n] = 0.0; } #endif @@ -1189,7 +1189,7 @@ __device__ void Average_Cell_All_Fields(int i, int j, int k, int nx, int ny, int mz = conserved[grid_enum::momentum_z * ncells + idn]; P = (conserved[grid_enum::Energy * ncells + idn] - (0.5 / d) * (mx * mx + my * my + mz * mz)) * (gamma - 1.0); #ifdef SCALAR - for (int n = 0; n < NSCALARS; n++) { //NOLINT + for (int n = 0; n < NSCALARS; n++) { // NOLINT scalar[n] = conserved[grid_enum::scalar * ncells + idn]; } #endif @@ -1200,7 +1200,7 @@ __device__ void Average_Cell_All_Fields(int i, int j, int k, int nx, int ny, int vz_av += mz; P_av += P / (gamma - 1.0); #ifdef SCALAR - for (int n = 0; n < NSCALARS; n++) { //NOLINT + for (int n = 0; n < NSCALARS; n++) { // NOLINT scalar_av[n] += scalar[n]; } #endif @@ -1215,7 +1215,7 @@ __device__ void Average_Cell_All_Fields(int i, int j, int k, int nx, int ny, int vy_av = vy_av / d_av; vz_av = vz_av / d_av; #ifdef SCALAR - for (int n = 0; n < NSCALARS; n++) { //NOLINT + for (int n = 0; n < NSCALARS; n++) { // NOLINT scalar_av[n] = scalar_av[n] / d_av; } #endif From cbf98efa96a37658b530642e83e3ad5a72644b20 Mon Sep 17 00:00:00 2001 From: Evan Schneider Date: Mon, 11 Dec 2023 14:01:30 -0500 Subject: [PATCH 6/6] missed one --- src/hydro/hydro_cuda.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/hydro/hydro_cuda.cu b/src/hydro/hydro_cuda.cu index 0127344af..1f1d07521 100644 --- a/src/hydro/hydro_cuda.cu +++ b/src/hydro/hydro_cuda.cu @@ -1232,7 +1232,7 @@ __device__ void Average_Cell_All_Fields(int i, int j, int k, int nx, int ny, int conserved[id + ncells * grid_enum::GasEnergy] = P_av / (gamma - 1.0); #endif #ifdef SCALAR - for (int n = 0; n < NSCALARS; n++) { + for (int n = 0; n < NSCALARS; n++) { // NOLINT conserved[id + ncells * grid_enum::scalar] = d_av * scalar_av[n]; } #endif