Skip to content

Commit

Permalink
widen clang format columns to 120
Browse files Browse the repository at this point in the history
  • Loading branch information
@alwinm
alwinm committed Feb 9, 2023
1 parent 4408600 commit 8e5b461
Show file tree
Hide file tree
Showing 158 changed files with 5,913 additions and 10,307 deletions.
2 changes: 1 addition & 1 deletion .clang-format
View file
Original file line number Diff line number Diff line change
Expand Up @@ -77,7 +77,7 @@ BreakConstructorInitializersBeforeComma: false
BreakConstructorInitializers: BeforeColon
BreakAfterJavaFieldAnnotations: false
BreakStringLiterals: true
ColumnLimit: 80
ColumnLimit: 120
CommentPragmas: '^ IWYU pragma:'
QualifierAlignment: Leave
CompactNamespaces: false
Expand Down
20 changes: 7 additions & 13 deletions src/analysis/analysis.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -84,8 +84,7 @@ void Grid3D::Compute_and_Output_Analysis(struct parameters *P)
chprintf("\nComputing Analysis \n");
#endif

cudaMemcpy(C.density, C.device, H.n_fields * H.n_cells * sizeof(Real),
cudaMemcpyDeviceToHost);
cudaMemcpy(C.density, C.device, H.n_fields * H.n_cells * sizeof(Real), cudaMemcpyDeviceToHost);

#ifdef PHASE_DIAGRAM
#ifdef CHEMISTRY_GPU
Expand All @@ -108,8 +107,7 @@ void Grid3D::Compute_and_Output_Analysis(struct parameters *P)
#endif

#ifdef LYA_STATISTICS
if (Analysis.Computed_Flux_Power_Spectrum == 1)
Analysis.Clear_Power_Spectrum_Measurements();
if (Analysis.Computed_Flux_Power_Spectrum == 1) Analysis.Clear_Power_Spectrum_Measurements();
#endif

#ifdef COSMOLOGY
Expand Down Expand Up @@ -139,17 +137,13 @@ void Grid3D::Initialize_Analysis_Module(struct parameters *P)
z_now = 0;
#endif

Analysis.Initialize(H.xdglobal, H.ydglobal, H.zdglobal, H.xblocal, H.yblocal,
H.zblocal, P->nx, P->ny, P->nz, H.nx_real, H.ny_real,
H.nz_real, H.dx, H.dy, H.dz, H.n_ghost, z_now, P);
Analysis.Initialize(H.xdglobal, H.ydglobal, H.zdglobal, H.xblocal, H.yblocal, H.zblocal, P->nx, P->ny, P->nz,
H.nx_real, H.ny_real, H.nz_real, H.dx, H.dy, H.dz, H.n_ghost, z_now, P);
}

void Analysis_Module::Initialize(Real Lx, Real Ly, Real Lz, Real x_min,
Real y_min, Real z_min, int nx, int ny, int nz,
int nx_real, int ny_real, int nz_real,
Real dx_real, Real dy_real, Real dz_real,
int n_ghost_hydro, Real z_now,
struct parameters *P)
void Analysis_Module::Initialize(Real Lx, Real Ly, Real Lz, Real x_min, Real y_min, Real z_min, int nx, int ny, int nz,
int nx_real, int ny_real, int nz_real, Real dx_real, Real dy_real, Real dz_real,
int n_ghost_hydro, Real z_now, struct parameters *P)
{
// Domain Length
Lbox_x = Lx;
Expand Down
7 changes: 3 additions & 4 deletions src/analysis/analysis.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -291,10 +291,9 @@ class Analysis_Module
#endif

Analysis_Module(void);
void Initialize(Real Lx, Real Ly, Real Lz, Real x_min, Real y_min, Real z_min,
int nx, int ny, int nz, int nx_real, int ny_real, int nz_real,
Real dx_real, Real dy_real, Real dz_real, int n_ghost_hydro,
Real z_now, struct parameters *P);
void Initialize(Real Lx, Real Ly, Real Lz, Real x_min, Real y_min, Real z_min, int nx, int ny, int nz, int nx_real,
int ny_real, int nz_real, Real dx_real, Real dy_real, Real dz_real, int n_ghost_hydro, Real z_now,
struct parameters *P);
void Reset(void);

void Load_Scale_Outputs(struct parameters *P);
Expand Down
48 changes: 17 additions & 31 deletions src/analysis/feedback_analysis.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -40,10 +40,8 @@ FeedbackAnalysis::FeedbackAnalysis(Grid3D& G)
}

#ifdef PARTICLES_GPU
CHECK(cudaMemcpy(d_circ_vel_x, h_circ_vel_x, G.H.n_cells * sizeof(Real),
cudaMemcpyHostToDevice));
CHECK(cudaMemcpy(d_circ_vel_y, h_circ_vel_y, G.H.n_cells * sizeof(Real),
cudaMemcpyHostToDevice));
CHECK(cudaMemcpy(d_circ_vel_x, h_circ_vel_x, G.H.n_cells * sizeof(Real), cudaMemcpyHostToDevice));
CHECK(cudaMemcpy(d_circ_vel_y, h_circ_vel_y, G.H.n_cells * sizeof(Real), cudaMemcpyHostToDevice));
#endif
}

Expand All @@ -69,10 +67,8 @@ void FeedbackAnalysis::Compute_Gas_Velocity_Dispersion(Grid3D& G)
Real x, y, z, r, xpm, xpp, ypm, ypp, zpm, zpp;
Real Pm, Pp;
Real dPdx, dPdy, dPdr;
Real vx, vy, vz, vrms_poisson, vrms_analytic, vcp, vca, vcxp, vcyp, vcxa,
vcya;
Real total_mass, partial_mass = 0, total_var_analytic = 0,
total_var_poisson = 0, partial_var_poisson = 0,
Real vx, vy, vz, vrms_poisson, vrms_analytic, vcp, vca, vcxp, vcyp, vcxa, vcya;
Real total_mass, partial_mass = 0, total_var_analytic = 0, total_var_poisson = 0, partial_var_poisson = 0,
partial_var_analytic = 0;

int n_ghost_grav = G.Particles.G.n_ghost_particles_grid;
Expand All @@ -83,8 +79,7 @@ void FeedbackAnalysis::Compute_Gas_Velocity_Dispersion(Grid3D& G)
for (k = 0; k < G.H.nz_real; k++) {
for (j = 0; j < G.H.ny_real; j++) {
for (i = 0; i < G.H.nx_real; i++) {
id = (i + G.H.n_ghost) + (j + G.H.n_ghost) * G.H.nx +
(k + G.H.n_ghost) * G.H.nx * G.H.ny;
id = (i + G.H.n_ghost) + (j + G.H.n_ghost) * G.H.nx + (k + G.H.n_ghost) * G.H.nx * G.H.ny;
partial_mass += G.C.density[id];
}
}
Expand All @@ -99,53 +94,44 @@ void FeedbackAnalysis::Compute_Gas_Velocity_Dispersion(Grid3D& G)
for (j = G.H.n_ghost; j < G.H.ny - G.H.n_ghost; j++) {
for (i = G.H.n_ghost; i < G.H.nx - G.H.n_ghost; i++) {
id = i + j * G.H.nx + k * G.H.nx * G.H.ny;
id_grav = (i + ghost_diff) + (j + ghost_diff) * nx_grav +
(k + ghost_diff) * nx_grav * ny_grav;
id_grav = (i + ghost_diff) + (j + ghost_diff) * nx_grav + (k + ghost_diff) * nx_grav * ny_grav;

if (G.C.density[id] < VRMS_CUTOFF_DENSITY)
continue; // in cgs, this is 0.01 cm^{-3}
if (G.C.density[id] < VRMS_CUTOFF_DENSITY) continue; // in cgs, this is 0.01 cm^{-3}

G.Get_Position(i, j, k, &x, &y, &z);
r = sqrt(x * x + y * y);

vcp = sqrt(r * fabs(G.Particles.G.gravity_x[id_grav] * x / r +
G.Particles.G.gravity_y[id_grav] * y / r));
vcp = sqrt(r * fabs(G.Particles.G.gravity_x[id_grav] * x / r + G.Particles.G.gravity_y[id_grav] * y / r));
vcxp = -y / r * vcp;
vcyp = x / r * vcp;
vx = G.C.momentum_x[id] / G.C.density[id];
vy = G.C.momentum_y[id] / G.C.density[id];
vz = G.C.momentum_z[id] / G.C.density[id];

partial_var_poisson +=
((vx - vcxp) * (vx - vcxp) + (vy - vcyp) * (vy - vcyp) + vz * vz) *
G.C.density[id];
partial_var_analytic +=
((vx - h_circ_vel_x[id]) * (vx - h_circ_vel_x[id]) +
(vy - h_circ_vel_y[id]) * (vy - h_circ_vel_y[id]) + (vz * vz)) *
G.C.density[id];
partial_var_poisson += ((vx - vcxp) * (vx - vcxp) + (vy - vcyp) * (vy - vcyp) + vz * vz) * G.C.density[id];
partial_var_analytic += ((vx - h_circ_vel_x[id]) * (vx - h_circ_vel_x[id]) +
(vy - h_circ_vel_y[id]) * (vy - h_circ_vel_y[id]) + (vz * vz)) *
G.C.density[id];
}
}
}
partial_var_poisson /= total_mass;
partial_var_analytic /= total_mass;

#ifdef MPI_CHOLLA
MPI_Reduce(&partial_var_poisson, &total_var_poisson, 1, MPI_CHREAL, MPI_SUM,
root, world);
MPI_Reduce(&partial_var_analytic, &total_var_analytic, 1, MPI_CHREAL, MPI_SUM,
root, world);
MPI_Reduce(&partial_var_poisson, &total_var_poisson, 1, MPI_CHREAL, MPI_SUM, root, world);
MPI_Reduce(&partial_var_analytic, &total_var_analytic, 1, MPI_CHREAL, MPI_SUM, root, world);

#else
total_var_poisson = partial_var_poisson;
total_var_analytic = partial_var_analytic;
#endif

vrms_poisson =
sqrt(total_var_poisson) * VELOCITY_UNIT / 1e5; // output in km/s
vrms_poisson = sqrt(total_var_poisson) * VELOCITY_UNIT / 1e5; // output in km/s
vrms_analytic = sqrt(total_var_analytic) * VELOCITY_UNIT / 1e5;

chprintf("feedback: time %f, dt=%f, vrms_p = %f km/s, vrms_a = %f km/s\n",
G.H.t, G.H.dt, vrms_poisson, vrms_analytic);
chprintf("feedback: time %f, dt=%f, vrms_p = %f km/s, vrms_a = %f km/s\n", G.H.t, G.H.dt, vrms_poisson,
vrms_analytic);

#elif defined(PARTICLES_GPU)
Compute_Gas_Velocity_Dispersion_GPU(G);
Expand Down
47 changes: 18 additions & 29 deletions src/analysis/feedback_analysis_gpu.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -8,9 +8,7 @@

#define MU 0.6
// in cgs, this is 0.01 cm^{-3}
#define MIN_DENSITY \
0.01 * MP *MU *LENGTH_UNIT *LENGTH_UNIT *LENGTH_UNIT / \
MASS_UNIT // 148279.7
#define MIN_DENSITY 0.01 * MP *MU *LENGTH_UNIT *LENGTH_UNIT *LENGTH_UNIT / MASS_UNIT // 148279.7
#define TPB_ANALYSIS 1024

__device__ void warpReduce(volatile Real *buff, size_t tid)
Expand All @@ -23,10 +21,8 @@ __device__ void warpReduce(volatile Real *buff, size_t tid)
if (TPB_ANALYSIS >= 2) buff[tid] += buff[tid + 1];
}

void __global__ Reduce_Tubulence_kernel(int nx, int ny, int nz, int n_ghost,
Real *density, Real *momentum_x,
Real *momentum_y, Real *momentum_z,
Real *circ_vel_x, Real *circ_vel_y,
void __global__ Reduce_Tubulence_kernel(int nx, int ny, int nz, int n_ghost, Real *density, Real *momentum_x,
Real *momentum_y, Real *momentum_z, Real *circ_vel_x, Real *circ_vel_y,
Real *partial_mass, Real *partial_vel)
{
__shared__ Real s_mass[TPB_ANALYSIS];
Expand All @@ -42,16 +38,15 @@ void __global__ Reduce_Tubulence_kernel(int nx, int ny, int nz, int n_ghost,
s_mass[tid] = 0;
s_vel[tid] = 0;
Real vx, vy, vz;
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[id] > MIN_DENSITY) {
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[id] > MIN_DENSITY) {
s_mass[tid] = density[id];
vx = momentum_x[id] / density[id];
vy = momentum_y[id] / density[id];
vz = momentum_z[id] / density[id];
s_vel[tid] = ((vx - circ_vel_x[id]) * (vx - circ_vel_x[id]) +
(vy - circ_vel_y[id]) * (vy - circ_vel_y[id]) + (vz * vz)) *
density[id];
s_vel[tid] =
((vx - circ_vel_x[id]) * (vx - circ_vel_x[id]) + (vy - circ_vel_y[id]) * (vy - circ_vel_y[id]) + (vz * vz)) *
density[id];
}
__syncthreads();

Expand All @@ -70,8 +65,7 @@ void __global__ Reduce_Tubulence_kernel(int nx, int ny, int nz, int n_ghost,
}
}

void __global__ Reduce_Tubulence_kernel_2(Real *input_m, Real *input_v,
Real *output_m, Real *output_v, int n)
void __global__ Reduce_Tubulence_kernel_2(Real *input_m, Real *input_v, Real *output_m, Real *output_v, int n)
{
__shared__ Real s_mass[TPB_ANALYSIS];
__shared__ Real s_vel[TPB_ANALYSIS];
Expand Down Expand Up @@ -147,36 +141,32 @@ void FeedbackAnalysis::Compute_Gas_Velocity_Dispersion_GPU(Grid3D &G)
Real total_mass = 0;
Real total_vel = 0;

hipLaunchKernelGGL(Reduce_Tubulence_kernel, ngrid, TPB_ANALYSIS, 0, 0, G.H.nx,
G.H.ny, G.H.nz, G.H.n_ghost, G.C.d_density,
G.C.d_momentum_x, G.C.d_momentum_y, G.C.d_momentum_z,
d_circ_vel_x, d_circ_vel_y, d_partial_mass, d_partial_vel);
hipLaunchKernelGGL(Reduce_Tubulence_kernel, ngrid, TPB_ANALYSIS, 0, 0, G.H.nx, G.H.ny, G.H.nz, G.H.n_ghost,
G.C.d_density, G.C.d_momentum_x, G.C.d_momentum_y, G.C.d_momentum_z, d_circ_vel_x, d_circ_vel_y,
d_partial_mass, d_partial_vel);

size_t n = ngrid;
Real *mass_input = d_partial_mass;
Real *vel_input = d_partial_vel;
while (n > TPB_ANALYSIS) {
ngrid = std::ceil((n * 1.) / TPB_ANALYSIS);
// printf("Reduce_Tubulence: Next kernel call grid size is %d\n", ngrid);
hipLaunchKernelGGL(Reduce_Tubulence_kernel_2, ngrid, TPB_ANALYSIS, 0, 0,
mass_input, vel_input, d_partial_mass, d_partial_vel, n);
hipLaunchKernelGGL(Reduce_Tubulence_kernel_2, ngrid, TPB_ANALYSIS, 0, 0, mass_input, vel_input, d_partial_mass,
d_partial_vel, n);
mass_input = d_partial_mass;
vel_input = d_partial_vel;
n = ngrid;
}

if (n > 1) {
hipLaunchKernelGGL(Reduce_Tubulence_kernel_2, 1, TPB_ANALYSIS, 0, 0,
d_partial_mass, d_partial_vel, d_partial_mass,
hipLaunchKernelGGL(Reduce_Tubulence_kernel_2, 1, TPB_ANALYSIS, 0, 0, d_partial_mass, d_partial_vel, d_partial_mass,
d_partial_vel, n);
}

// cudaDeviceSynchronize();

CHECK(cudaMemcpy(h_partial_mass, d_partial_mass, ngrid * sizeof(Real),
cudaMemcpyDeviceToHost));
CHECK(cudaMemcpy(h_partial_vel, d_partial_vel, ngrid * sizeof(Real),
cudaMemcpyDeviceToHost));
CHECK(cudaMemcpy(h_partial_mass, d_partial_mass, ngrid * sizeof(Real), cudaMemcpyDeviceToHost));
CHECK(cudaMemcpy(h_partial_vel, d_partial_vel, ngrid * sizeof(Real), cudaMemcpyDeviceToHost));

#ifdef MPI_CHOLLA
MPI_Allreduce(h_partial_mass, &total_mass, 1, MPI_CHREAL, MPI_SUM, world);
Expand All @@ -187,8 +177,7 @@ void FeedbackAnalysis::Compute_Gas_Velocity_Dispersion_GPU(Grid3D &G)
#endif

if (total_vel < 0 || total_mass < 0) {
chprintf("feedback trouble. total_vel = %.3e, total_mass = %.3e\n",
total_vel, total_mass);
chprintf("feedback trouble. total_vel = %.3e, total_mass = %.3e\n", total_vel, total_mass);
}

chprintf("feedback: time %f, dt=%f, vrms = %f km/s\n", G.H.t, G.H.dt,
Expand Down
Loading

0 comments on commit 8e5b461

Please sign in to comment.