From acd2108b93039c9f598c3d3fd770c337afdff196 Mon Sep 17 00:00:00 2001 From: nlg550 <38725499+nlg550@users.noreply.github.com> Date: Fri, 9 Apr 2021 16:02:03 -0300 Subject: [PATCH] Added weak scaling tests --- parallel/ompss2_openacc/Makefile | 4 +- .../input/weak/cold-2000-1073M-4096-4096.c | 57 +++++++++++++++++ .../input/weak/cold-2000-268M-2048-2048.c | 57 +++++++++++++++++ .../input/weak/cold-2000-538M-2900-2900.c | 56 +++++++++++++++++ .../input/weak/cold-2000-829M-3600-3600.c | 56 +++++++++++++++++ .../input/weak/warm-2000-1073M-4096-4096.c | 57 +++++++++++++++++ .../input/weak/warm-2000-268M-2048-2048.c | 57 +++++++++++++++++ .../input/weak/warm-2000-538M-2900-2900.c | 57 +++++++++++++++++ .../input/weak/warm-2000-829M-3600-3600.c | 57 +++++++++++++++++ .../input/weak/weibel-2000-151M-2048-2048.c | 61 +++++++++++++++++++ .../input/weak/weibel-2000-303M-2900-2900.c | 61 +++++++++++++++++++ .../input/weak/weibel-2000-467M-3600-3600.c | 61 +++++++++++++++++++ .../input/weak/weibel-2000-604M-4096-4096.c | 61 +++++++++++++++++++ parallel/ompss2_openacc/kernel_particles.c | 11 ++++ parallel/ompss2_openacc/main.c | 21 ++++++- .../ompss2_openacc/set-ompss-openacc-env.sh | 6 ++ parallel/ompss2_openacc/simulation.c | 9 ++- 17 files changed, 743 insertions(+), 6 deletions(-) create mode 100644 parallel/ompss2_openacc/input/weak/cold-2000-1073M-4096-4096.c create mode 100644 parallel/ompss2_openacc/input/weak/cold-2000-268M-2048-2048.c create mode 100644 parallel/ompss2_openacc/input/weak/cold-2000-538M-2900-2900.c create mode 100644 parallel/ompss2_openacc/input/weak/cold-2000-829M-3600-3600.c create mode 100644 parallel/ompss2_openacc/input/weak/warm-2000-1073M-4096-4096.c create mode 100644 parallel/ompss2_openacc/input/weak/warm-2000-268M-2048-2048.c create mode 100644 parallel/ompss2_openacc/input/weak/warm-2000-538M-2900-2900.c create mode 100644 parallel/ompss2_openacc/input/weak/warm-2000-829M-3600-3600.c create mode 100644 parallel/ompss2_openacc/input/weak/weibel-2000-151M-2048-2048.c create mode 100644 parallel/ompss2_openacc/input/weak/weibel-2000-303M-2900-2900.c create mode 100644 parallel/ompss2_openacc/input/weak/weibel-2000-467M-3600-3600.c create mode 100644 parallel/ompss2_openacc/input/weak/weibel-2000-604M-4096-4096.c create mode 100644 parallel/ompss2_openacc/set-ompss-openacc-env.sh diff --git a/parallel/ompss2_openacc/Makefile b/parallel/ompss2_openacc/Makefile index 1f07add..69a190d 100644 --- a/parallel/ompss2_openacc/Makefile +++ b/parallel/ompss2_openacc/Makefile @@ -1,9 +1,9 @@ # GCC options CC = pgimcc -CFLAGS = --ompss-2 -O3 --openacc -DTEST --cuda +CFLAGS = --ompss-2 -O3 --openacc --cuda #CC = pgcc -#CFLAGS = -acc -ta=tesla:managed -Minfo=acc -O3 -fast -cudalibs -Mcuda=ptxinfo -g +#CFLAGS = -acc -ta=tesla:managed:cuda10.1 -Minfo=acc -O3 -fast -Mcuda=ptxinfo -g INCLUDES = LDFLAGS = -lm -lcuda diff --git a/parallel/ompss2_openacc/input/weak/cold-2000-1073M-4096-4096.c b/parallel/ompss2_openacc/input/weak/cold-2000-1073M-4096-4096.c new file mode 100644 index 0000000..a767997 --- /dev/null +++ b/parallel/ompss2_openacc/input/weak/cold-2000-1073M-4096-4096.c @@ -0,0 +1,57 @@ +/** + * ZPIC - em2d + * + * Weibel instability + */ + +#include +#include "../../simulation.h" + +void sim_init(t_simulation *sim, int n_regions, float gpu_percentage, int n_gpu_regions) +{ + // Time step + float dt = 0.035; + float tmax = 70.0; + + // Simulation box + int nx[2] = {4096, 4096}; + float box[2] = {204.8, 204.8}; + + // Diagnostic frequency + int ndump = 500; + + // Initialize particles + const int n_species = 1; + t_species *species = (t_species*) malloc(n_species * sizeof(t_species)); + + // Use 8x8 particles per cell + int ppc[] = {8, 8}; + + // Initial fluid and thermal velocities + t_part_data ufl[] = {0.0, 0.0, 0.0}; + t_part_data uth[] = {0.0, 0.0, 0.0}; + + spec_new(&species[0], "electrons", -1.0, ppc, ufl, uth, nx, box, dt, NULL, nx[1], -1); + + // Initialize Simulation data + sim_new(sim, nx, box, dt, tmax, ndump, species, n_species, "cold-2000-1073M-4096-4096", n_regions, gpu_percentage, n_gpu_regions); + + free(species); +} + +void sim_report(t_simulation *sim) +{ + //sim_report_csv(sim); + sim_report_energy(sim); + + // Bx, By, Bz + sim_report_grid_zdf(sim, REPORT_BFLD, 0); + sim_report_grid_zdf(sim, REPORT_BFLD, 1); + sim_report_grid_zdf(sim, REPORT_BFLD, 2); + + // Jz + sim_report_grid_zdf(sim, REPORT_CURRENT, 2); + + // electron and positron density + sim_report_spec_zdf(sim, 0, CHARGE, NULL, NULL); +} diff --git a/parallel/ompss2_openacc/input/weak/cold-2000-268M-2048-2048.c b/parallel/ompss2_openacc/input/weak/cold-2000-268M-2048-2048.c new file mode 100644 index 0000000..1d1e445 --- /dev/null +++ b/parallel/ompss2_openacc/input/weak/cold-2000-268M-2048-2048.c @@ -0,0 +1,57 @@ +/** + * ZPIC - em2d + * + * Weibel instability + */ + +#include +#include "../../simulation.h" + +void sim_init(t_simulation *sim, int n_regions, float gpu_percentage, int n_gpu_regions) +{ + // Time step + float dt = 0.035; + float tmax = 70.0; + + // Simulation box + int nx[2] = {2048, 2048}; + float box[2] = {102.4, 102.4}; + + // Diagnostic frequency + int ndump = 500; + + // Initialize particles + const int n_species = 1; + t_species *species = (t_species*) malloc(n_species * sizeof(t_species)); + + // Use 8x8 particles per cell + int ppc[] = {8, 8}; + + // Initial fluid and thermal velocities + t_part_data ufl[] = {0.0, 0.0, 0.0}; + t_part_data uth[] = {0.0, 0.0, 0.0}; + + spec_new(&species[0], "electrons", -1.0, ppc, ufl, uth, nx, box, dt, NULL, nx[1], -1); + + // Initialize Simulation data + sim_new(sim, nx, box, dt, tmax, ndump, species, n_species, "cold-2000-268M-2048-2048", n_regions, gpu_percentage, n_gpu_regions); + + free(species); +} + +void sim_report(t_simulation *sim) +{ + //sim_report_csv(sim); + sim_report_energy(sim); + + // Bx, By, Bz + sim_report_grid_zdf(sim, REPORT_BFLD, 0); + sim_report_grid_zdf(sim, REPORT_BFLD, 1); + sim_report_grid_zdf(sim, REPORT_BFLD, 2); + + // Jz + sim_report_grid_zdf(sim, REPORT_CURRENT, 2); + + // electron and positron density + sim_report_spec_zdf(sim, 0, CHARGE, NULL, NULL); +} diff --git a/parallel/ompss2_openacc/input/weak/cold-2000-538M-2900-2900.c b/parallel/ompss2_openacc/input/weak/cold-2000-538M-2900-2900.c new file mode 100644 index 0000000..ecf7774 --- /dev/null +++ b/parallel/ompss2_openacc/input/weak/cold-2000-538M-2900-2900.c @@ -0,0 +1,56 @@ +/** + * ZPIC - em2d + * + * Weibel instability + */ + +#include +#include "../../simulation.h" + +void sim_init(t_simulation *sim, int n_regions, float gpu_percentage, int n_gpu_regions) +{ + // Time step + float dt = 0.035; + float tmax = 70.0; + + // Simulation box + int nx[2] = {2900, 2900}; + float box[2] = {145, 145}; + + // Diagnostic frequency + int ndump = 500; + + // Initialize particles + const int n_species = 1; + t_species *species = (t_species*) malloc(n_species * sizeof(t_species)); + + // Use 8x8 particles per cell + int ppc[] = {8, 8}; + + // Initial fluid and thermal velocities + t_part_data ufl[] = {0.0, 0.0, 0.0}; + t_part_data uth[] = {0.0, 0.0, 0.0}; + + spec_new(&species[0], "electrons", -1.0, ppc, ufl, uth, nx, box, dt, NULL, nx[1], -1); + + // Initialize Simulation data + sim_new(sim, nx, box, dt, tmax, ndump, species, n_species, "cold-2000-538M-2900-2900", n_regions, gpu_percentage, n_gpu_regions); + free(species); +} + +void sim_report(t_simulation *sim) +{ + //sim_report_csv(sim); + sim_report_energy(sim); + + // Bx, By, Bz + sim_report_grid_zdf(sim, REPORT_BFLD, 0); + sim_report_grid_zdf(sim, REPORT_BFLD, 1); + sim_report_grid_zdf(sim, REPORT_BFLD, 2); + + // Jz + sim_report_grid_zdf(sim, REPORT_CURRENT, 2); + + // electron and positron density + sim_report_spec_zdf(sim, 0, CHARGE, NULL, NULL); +} diff --git a/parallel/ompss2_openacc/input/weak/cold-2000-829M-3600-3600.c b/parallel/ompss2_openacc/input/weak/cold-2000-829M-3600-3600.c new file mode 100644 index 0000000..37e2324 --- /dev/null +++ b/parallel/ompss2_openacc/input/weak/cold-2000-829M-3600-3600.c @@ -0,0 +1,56 @@ +/** + * ZPIC - em2d + * + * Weibel instability + */ + +#include +#include "../../simulation.h" + +void sim_init(t_simulation *sim, int n_regions, float gpu_percentage, int n_gpu_regions) +{ + // Time step + float dt = 0.035; + float tmax = 70.0; + + // Simulation box + int nx[2] = {3600, 3600}; + float box[2] = {180, 180}; + + // Diagnostic frequency + int ndump = 500; + + // Initialize particles + const int n_species = 1; + t_species *species = (t_species*) malloc(n_species * sizeof(t_species)); + + // Use 8x8 particles per cell + int ppc[] = {8, 8}; + + // Initial fluid and thermal velocities + t_part_data ufl[] = {0.0, 0.0, 0.0}; + t_part_data uth[] = {0.0, 0.0, 0.0}; + + spec_new(&species[0], "electrons", -1.0, ppc, ufl, uth, nx, box, dt, NULL, nx[1], -1); + + // Initialize Simulation data + sim_new(sim, nx, box, dt, tmax, ndump, species, n_species, "cold-2000-829M-3600-3600", n_regions, gpu_percentage, n_gpu_regions); + free(species); +} + +void sim_report(t_simulation *sim) +{ + //sim_report_csv(sim); + sim_report_energy(sim); + + // Bx, By, Bz + sim_report_grid_zdf(sim, REPORT_BFLD, 0); + sim_report_grid_zdf(sim, REPORT_BFLD, 1); + sim_report_grid_zdf(sim, REPORT_BFLD, 2); + + // Jz + sim_report_grid_zdf(sim, REPORT_CURRENT, 2); + + // electron and positron density + sim_report_spec_zdf(sim, 0, CHARGE, NULL, NULL); +} diff --git a/parallel/ompss2_openacc/input/weak/warm-2000-1073M-4096-4096.c b/parallel/ompss2_openacc/input/weak/warm-2000-1073M-4096-4096.c new file mode 100644 index 0000000..f9fb84c --- /dev/null +++ b/parallel/ompss2_openacc/input/weak/warm-2000-1073M-4096-4096.c @@ -0,0 +1,57 @@ +/** + * ZPIC - em2d + * + * Weibel instability + */ + +#include +#include "../../simulation.h" + +void sim_init(t_simulation *sim, int n_regions, float gpu_percentage, int n_gpu_regions) +{ + // Time step + float dt = 0.035; + float tmax = 70.0; + + // Simulation box + int nx[2] = {4096, 4096}; + float box[2] = {204.8, 204.8}; + + // Diagnostic frequency + int ndump = 500; + + // Initialize particles + const int n_species = 1; + t_species *species = (t_species*) malloc(n_species * sizeof(t_species)); + + // Use 8x8 particles per cell + int ppc[] = {8, 8}; + + // Initial fluid and thermal velocities + t_part_data ufl[] = {0.0, 0.0, 0.0}; + t_part_data uth[] = {0.01, 0.01, 0.01}; + + spec_new(&species[0], "electrons", -1.0, ppc, ufl, uth, nx, box, dt, NULL, nx[1], -1); + + // Initialize Simulation data + sim_new(sim, nx, box, dt, tmax, ndump, species, n_species, "warm-2000-1073M-4096-4096", n_regions, gpu_percentage, n_gpu_regions); + + free(species); +} + +void sim_report(t_simulation *sim) +{ + //sim_report_csv(sim); + sim_report_energy(sim); + + // Bx, By, Bz + sim_report_grid_zdf(sim, REPORT_BFLD, 0); + sim_report_grid_zdf(sim, REPORT_BFLD, 1); + sim_report_grid_zdf(sim, REPORT_BFLD, 2); + + // Jz + sim_report_grid_zdf(sim, REPORT_CURRENT, 2); + + // electron and positron density + sim_report_spec_zdf(sim, 0, CHARGE, NULL, NULL); +} diff --git a/parallel/ompss2_openacc/input/weak/warm-2000-268M-2048-2048.c b/parallel/ompss2_openacc/input/weak/warm-2000-268M-2048-2048.c new file mode 100644 index 0000000..e4c63ce --- /dev/null +++ b/parallel/ompss2_openacc/input/weak/warm-2000-268M-2048-2048.c @@ -0,0 +1,57 @@ +/** + * ZPIC - em2d + * + * Weibel instability + */ + +#include +#include "../../simulation.h" + +void sim_init(t_simulation *sim, int n_regions, float gpu_percentage, int n_gpu_regions) +{ + // Time step + float dt = 0.035; + float tmax = 70.0; + + // Simulation box + int nx[2] = {2048, 2048}; + float box[2] = {102.4, 102.4}; + + // Diagnostic frequency + int ndump = 500; + + // Initialize particles + const int n_species = 1; + t_species *species = (t_species*) malloc(n_species * sizeof(t_species)); + + // Use 8x8 particles per cell + int ppc[] = {8, 8}; + + // Initial fluid and thermal velocities + t_part_data ufl[] = {0.0, 0.0, 0.0}; + t_part_data uth[] = {0.01, 0.01, 0.01}; + + spec_new(&species[0], "electrons", -1.0, ppc, ufl, uth, nx, box, dt, NULL, nx[1], -1); + + // Initialize Simulation data + sim_new(sim, nx, box, dt, tmax, ndump, species, n_species, "warm-2000-268M-2048-2048", n_regions, gpu_percentage, n_gpu_regions); + + free(species); +} + +void sim_report(t_simulation *sim) +{ + //sim_report_csv(sim); + sim_report_energy(sim); + + // Bx, By, Bz + sim_report_grid_zdf(sim, REPORT_BFLD, 0); + sim_report_grid_zdf(sim, REPORT_BFLD, 1); + sim_report_grid_zdf(sim, REPORT_BFLD, 2); + + // Jz + sim_report_grid_zdf(sim, REPORT_CURRENT, 2); + + // electron and positron density + sim_report_spec_zdf(sim, 0, CHARGE, NULL, NULL); +} diff --git a/parallel/ompss2_openacc/input/weak/warm-2000-538M-2900-2900.c b/parallel/ompss2_openacc/input/weak/warm-2000-538M-2900-2900.c new file mode 100644 index 0000000..40f61bc --- /dev/null +++ b/parallel/ompss2_openacc/input/weak/warm-2000-538M-2900-2900.c @@ -0,0 +1,57 @@ +/** + * ZPIC - em2d + * + * Weibel instability + */ + +#include +#include "../../simulation.h" + +void sim_init(t_simulation *sim, int n_regions, float gpu_percentage, int n_gpu_regions) +{ + // Time step + float dt = 0.035; + float tmax = 70.0; + + // Simulation box + int nx[2] = {2900, 2900}; + float box[2] = {145, 145}; + + // Diagnostic frequency + int ndump = 500; + + // Initialize particles + const int n_species = 1; + t_species *species = (t_species*) malloc(n_species * sizeof(t_species)); + + // Use 8x8 particles per cell + int ppc[] = {8, 8}; + + // Initial fluid and thermal velocities + t_part_data ufl[] = {0.0, 0.0, 0.0}; + t_part_data uth[] = {0.01, 0.01, 0.01}; + + spec_new(&species[0], "electrons", -1.0, ppc, ufl, uth, nx, box, dt, NULL, nx[1], -1); + + // Initialize Simulation data + sim_new(sim, nx, box, dt, tmax, ndump, species, n_species, "warm-2000-538M-2900-2900", n_regions, gpu_percentage, n_gpu_regions); + + free(species); +} + +void sim_report(t_simulation *sim) +{ + //sim_report_csv(sim); + sim_report_energy(sim); + + // Bx, By, Bz + sim_report_grid_zdf(sim, REPORT_BFLD, 0); + sim_report_grid_zdf(sim, REPORT_BFLD, 1); + sim_report_grid_zdf(sim, REPORT_BFLD, 2); + + // Jz + sim_report_grid_zdf(sim, REPORT_CURRENT, 2); + + // electron and positron density + sim_report_spec_zdf(sim, 0, CHARGE, NULL, NULL); +} diff --git a/parallel/ompss2_openacc/input/weak/warm-2000-829M-3600-3600.c b/parallel/ompss2_openacc/input/weak/warm-2000-829M-3600-3600.c new file mode 100644 index 0000000..161a80d --- /dev/null +++ b/parallel/ompss2_openacc/input/weak/warm-2000-829M-3600-3600.c @@ -0,0 +1,57 @@ +/** + * ZPIC - em2d + * + * Weibel instability + */ + +#include +#include "../../simulation.h" + +void sim_init(t_simulation *sim, int n_regions, float gpu_percentage, int n_gpu_regions) +{ + // Time step + float dt = 0.035; + float tmax = 70.0; + + // Simulation box + int nx[2] = {3600, 3600}; + float box[2] = {180, 180}; + + // Diagnostic frequency + int ndump = 500; + + // Initialize particles + const int n_species = 1; + t_species *species = (t_species*) malloc(n_species * sizeof(t_species)); + + // Use 8x8 particles per cell + int ppc[] = {8, 8}; + + // Initial fluid and thermal velocities + t_part_data ufl[] = {0.0, 0.0, 0.0}; + t_part_data uth[] = {0.01, 0.01, 0.01}; + + spec_new(&species[0], "electrons", -1.0, ppc, ufl, uth, nx, box, dt, NULL, nx[1], -1); + + // Initialize Simulation data + sim_new(sim, nx, box, dt, tmax, ndump, species, n_species, "warm-2000-829M-3600-3600", n_regions, gpu_percentage, n_gpu_regions); + + free(species); +} + +void sim_report(t_simulation *sim) +{ + //sim_report_csv(sim); + sim_report_energy(sim); + + // Bx, By, Bz + sim_report_grid_zdf(sim, REPORT_BFLD, 0); + sim_report_grid_zdf(sim, REPORT_BFLD, 1); + sim_report_grid_zdf(sim, REPORT_BFLD, 2); + + // Jz + sim_report_grid_zdf(sim, REPORT_CURRENT, 2); + + // electron and positron density + sim_report_spec_zdf(sim, 0, CHARGE, NULL, NULL); +} diff --git a/parallel/ompss2_openacc/input/weak/weibel-2000-151M-2048-2048.c b/parallel/ompss2_openacc/input/weak/weibel-2000-151M-2048-2048.c new file mode 100644 index 0000000..52d1397 --- /dev/null +++ b/parallel/ompss2_openacc/input/weak/weibel-2000-151M-2048-2048.c @@ -0,0 +1,61 @@ +/** + * ZPIC - em2d + * + * Weibel instability + */ + +#include +#include "../../simulation.h" + +void sim_init(t_simulation *sim, int n_regions, float gpu_percentage, int n_gpu_regions) +{ + // Time step + float dt = 0.035; + float tmax = 70.0; + + // Simulation box + int nx[2] = {2048, 2048}; + float box[2] = {102.4, 102.4}; + + // Diagnostic frequency + int ndump = 500; + + // Initialize particles + const int n_species = 2; + t_species *species = (t_species*) malloc(n_species * sizeof(t_species)); + + // Use 8x8 particles per cell + int ppc[] = {6, 6}; + + // Initial fluid and thermal velocities + t_part_data ufl[] = {0.0, 0.0, 0.6}; + t_part_data uth[] = {0.1, 0.1, 0.1}; + + spec_new(&species[0], "electrons", -1.0, ppc, ufl, uth, nx, box, dt, NULL, nx[1], -1); + + ufl[2] = -ufl[2]; + spec_new(&species[1], "positrons", +1.0, ppc, ufl, uth, nx, box, dt, NULL, nx[1], -1); + + // Initialize Simulation data + sim_new(sim, nx, box, dt, tmax, ndump, species, n_species, "weibel-2000-151M-2048-2048", n_regions, gpu_percentage, n_gpu_regions); + + free(species); +} + +void sim_report(t_simulation *sim) +{ + //sim_report_csv(sim); + sim_report_energy(sim); + + // Bx, By, Bz + sim_report_grid_zdf(sim, REPORT_BFLD, 0); + sim_report_grid_zdf(sim, REPORT_BFLD, 1); + sim_report_grid_zdf(sim, REPORT_BFLD, 2); + + // Jz + sim_report_grid_zdf(sim, REPORT_CURRENT, 2); + + // electron and positron density + sim_report_spec_zdf(sim, 0, CHARGE, NULL, NULL); + sim_report_spec_zdf(sim, 1, CHARGE, NULL, NULL); +} diff --git a/parallel/ompss2_openacc/input/weak/weibel-2000-303M-2900-2900.c b/parallel/ompss2_openacc/input/weak/weibel-2000-303M-2900-2900.c new file mode 100644 index 0000000..40466a3 --- /dev/null +++ b/parallel/ompss2_openacc/input/weak/weibel-2000-303M-2900-2900.c @@ -0,0 +1,61 @@ +/** + * ZPIC - em2d + * + * Weibel instability + */ + +#include +#include "../../simulation.h" + +void sim_init(t_simulation *sim, int n_regions, float gpu_percentage, int n_gpu_regions) +{ + // Time step + float dt = 0.035; + float tmax = 70.0; + + // Simulation box + int nx[2] = {2900, 2900}; + float box[2] = {145, 145}; + + // Diagnostic frequency + int ndump = 500; + + // Initialize particles + const int n_species = 2; + t_species *species = (t_species*) malloc(n_species * sizeof(t_species)); + + // Use 8x8 particles per cell + int ppc[] = {6, 6}; + + // Initial fluid and thermal velocities + t_part_data ufl[] = {0.0, 0.0, 0.6}; + t_part_data uth[] = {0.1, 0.1, 0.1}; + + spec_new(&species[0], "electrons", -1.0, ppc, ufl, uth, nx, box, dt, NULL, nx[1], -1); + + ufl[2] = -ufl[2]; + spec_new(&species[1], "positrons", +1.0, ppc, ufl, uth, nx, box, dt, NULL, nx[1], -1); + + // Initialize Simulation data + sim_new(sim, nx, box, dt, tmax, ndump, species, n_species, "weibel-2000-303M-2900-2900", n_regions, gpu_percentage, n_gpu_regions); + + free(species); +} + +void sim_report(t_simulation *sim) +{ + //sim_report_csv(sim); + sim_report_energy(sim); + + // Bx, By, Bz + sim_report_grid_zdf(sim, REPORT_BFLD, 0); + sim_report_grid_zdf(sim, REPORT_BFLD, 1); + sim_report_grid_zdf(sim, REPORT_BFLD, 2); + + // Jz + sim_report_grid_zdf(sim, REPORT_CURRENT, 2); + + // electron and positron density + sim_report_spec_zdf(sim, 0, CHARGE, NULL, NULL); + sim_report_spec_zdf(sim, 1, CHARGE, NULL, NULL); +} diff --git a/parallel/ompss2_openacc/input/weak/weibel-2000-467M-3600-3600.c b/parallel/ompss2_openacc/input/weak/weibel-2000-467M-3600-3600.c new file mode 100644 index 0000000..44b9a8e --- /dev/null +++ b/parallel/ompss2_openacc/input/weak/weibel-2000-467M-3600-3600.c @@ -0,0 +1,61 @@ +/** + * ZPIC - em2d + * + * Weibel instability + */ + +#include +#include "../../simulation.h" + +void sim_init(t_simulation *sim, int n_regions, float gpu_percentage, int n_gpu_regions) +{ + // Time step + float dt = 0.035; + float tmax = 70.0; + + // Simulation box + int nx[2] = {3600, 3600}; + float box[2] = {180, 180}; + + // Diagnostic frequency + int ndump = 500; + + // Initialize particles + const int n_species = 2; + t_species *species = (t_species*) malloc(n_species * sizeof(t_species)); + + // Use 8x8 particles per cell + int ppc[] = {6, 6}; + + // Initial fluid and thermal velocities + t_part_data ufl[] = {0.0, 0.0, 0.6}; + t_part_data uth[] = {0.1, 0.1, 0.1}; + + spec_new(&species[0], "electrons", -1.0, ppc, ufl, uth, nx, box, dt, NULL, nx[1], -1); + + ufl[2] = -ufl[2]; + spec_new(&species[1], "positrons", +1.0, ppc, ufl, uth, nx, box, dt, NULL, nx[1], -1); + + // Initialize Simulation data + sim_new(sim, nx, box, dt, tmax, ndump, species, n_species, "weibel-2000-467M-3600-3600", n_regions, gpu_percentage, n_gpu_regions); + + free(species); +} + +void sim_report(t_simulation *sim) +{ + //sim_report_csv(sim); + sim_report_energy(sim); + + // Bx, By, Bz + sim_report_grid_zdf(sim, REPORT_BFLD, 0); + sim_report_grid_zdf(sim, REPORT_BFLD, 1); + sim_report_grid_zdf(sim, REPORT_BFLD, 2); + + // Jz + sim_report_grid_zdf(sim, REPORT_CURRENT, 2); + + // electron and positron density + sim_report_spec_zdf(sim, 0, CHARGE, NULL, NULL); + sim_report_spec_zdf(sim, 1, CHARGE, NULL, NULL); +} diff --git a/parallel/ompss2_openacc/input/weak/weibel-2000-604M-4096-4096.c b/parallel/ompss2_openacc/input/weak/weibel-2000-604M-4096-4096.c new file mode 100644 index 0000000..74c27f8 --- /dev/null +++ b/parallel/ompss2_openacc/input/weak/weibel-2000-604M-4096-4096.c @@ -0,0 +1,61 @@ +/** + * ZPIC - em2d + * + * Weibel instability + */ + +#include +#include "../../simulation.h" + +void sim_init(t_simulation *sim, int n_regions, float gpu_percentage, int n_gpu_regions) +{ + // Time step + float dt = 0.035; + float tmax = 70.0; + + // Simulation box + int nx[2] = {4096, 4096}; + float box[2] = {204.8, 204.8}; + + // Diagnostic frequency + int ndump = 500; + + // Initialize particles + const int n_species = 2; + t_species *species = (t_species*) malloc(n_species * sizeof(t_species)); + + // Use 8x8 particles per cell + int ppc[] = {6, 6}; + + // Initial fluid and thermal velocities + t_part_data ufl[] = {0.0, 0.0, 0.6}; + t_part_data uth[] = {0.1, 0.1, 0.1}; + + spec_new(&species[0], "electrons", -1.0, ppc, ufl, uth, nx, box, dt, NULL, nx[1], -1); + + ufl[2] = -ufl[2]; + spec_new(&species[1], "positrons", +1.0, ppc, ufl, uth, nx, box, dt, NULL, nx[1], -1); + + // Initialize Simulation data + sim_new(sim, nx, box, dt, tmax, ndump, species, n_species, "weibel-2000-604M-4096-4096", n_regions, gpu_percentage, n_gpu_regions); + + free(species); +} + +void sim_report(t_simulation *sim) +{ + //sim_report_csv(sim); + sim_report_energy(sim); + + // Bx, By, Bz + sim_report_grid_zdf(sim, REPORT_BFLD, 0); + sim_report_grid_zdf(sim, REPORT_BFLD, 1); + sim_report_grid_zdf(sim, REPORT_BFLD, 2); + + // Jz + sim_report_grid_zdf(sim, REPORT_CURRENT, 2); + + // electron and positron density + sim_report_spec_zdf(sim, 0, CHARGE, NULL, NULL); + sim_report_spec_zdf(sim, 1, CHARGE, NULL, NULL); +} diff --git a/parallel/ompss2_openacc/kernel_particles.c b/parallel/ompss2_openacc/kernel_particles.c index 7e8317a..1e01976 100644 --- a/parallel/ompss2_openacc/kernel_particles.c +++ b/parallel/ompss2_openacc/kernel_particles.c @@ -213,8 +213,10 @@ void spec_move_vector_int_full(int *restrict vector, int *restrict new_pos, cons { int *restrict temp = malloc(size * sizeof(int)); + #pragma acc parallel loop for(int i = 0; i < size; i++) temp[i] = vector[i]; + #pragma acc parallel loop for(int i = 0; i < size; i++) if(new_pos[i] >= 0) vector[new_pos[i]] = temp[i]; @@ -226,8 +228,10 @@ void spec_move_vector_float_full(float *restrict vector, int *restrict new_pos, { float *restrict temp = malloc(size * sizeof(float)); + #pragma acc parallel loop for(int i = 0; i < size; i++) temp[i] = vector[i]; + #pragma acc parallel loop for(int i = 0; i < size; i++) if(new_pos[i] >= 0) vector[new_pos[i]] = temp[i]; @@ -254,17 +258,23 @@ void spec_organize_in_tiles(t_species *spec, const int limits_y[2], const int de memset(spec->mv_part_offset, 0, (n_tiles_y * n_tiles_x + 1) * sizeof(int)); // Calculate the histogram (number of particles per tile) + #pragma acc set device_num(device) + #pragma acc parallel loop for (int i = 0; i < size; i++) { int ix = spec->main_vector.ix[i] / TILE_SIZE; int iy = (spec->main_vector.iy[i] - limits_y[0]) / TILE_SIZE; + + #pragma acc atomic capture pos[i] = tile_offset[ix + iy * n_tiles_x]++; } + // Prefix sum to find the initial idx of each tile in the particle vector prefix_sum_serial(tile_offset, n_tiles_x * n_tiles_y + 1); // Calculate the target position of each particle + #pragma acc parallel loop for (int i = 0; i < size; i++) { int ix = spec->main_vector.ix[i] / TILE_SIZE; @@ -286,6 +296,7 @@ void spec_organize_in_tiles(t_species *spec, const int limits_y[2], const int de spec_move_vector_float_full(spec->main_vector.uz, pos, size); // Validate all the particles + #pragma acc parallel loop for (int k = 0; k < final_size; k++) spec->main_vector.invalid[k] = false; diff --git a/parallel/ompss2_openacc/main.c b/parallel/ompss2_openacc/main.c index b0ce0c5..c84668e 100644 --- a/parallel/ompss2_openacc/main.c +++ b/parallel/ompss2_openacc/main.c @@ -28,12 +28,27 @@ #include "timer.h" // Simulation parameters (naming scheme : ---.c) -#include "input/weibel-2000-151M-2048-2048.c" +//#include "input/weibel-2000-151M-2048-2048.c" // #include "input/lwfa-8000-131M-4000-2048.c" //#include "input/warm-2000-538M-2900-2900.c" -//#include "input/weibel-500-4M-512-512.c" +//#include "input/weibel-500-67M-512-512.c" -//#pragma oss assert("version.dependencies==regions") +//#include "input/weak/weibel-2000-151M-2048-2048.c" +//#include "input/weak/weibel-2000-303M-2900-2900.c" +//#include "input/weak/weibel-2000-467M-3600-3600.c" +//#include "input/weak/weibel-2000-604M-4096-4096.c" + +//#include "input/weak/warm-2000-268M-2048-2048.c" +//#include "input/weak/warm-2000-538M-2900-2900.c" +//#include "input/weak/warm-2000-829M-3600-3600.c" +#include "input/weak/warm-2000-1073M-4096-4096.c" + +//#include "input/weak/cold-2000-268M-2048-2048.c" +//#include "input/weak/cold-2000-538M-2900-2900.c" +//#include "input/weak/cold-2000-829M-3600-3600.c" +//#include "input/weak/cold-2000-1073M-4096-4096.c" + +#pragma oss assert("version.dependencies==regions") int main(int argc, const char *argv[]) { if(argc != 4) diff --git a/parallel/ompss2_openacc/set-ompss-openacc-env.sh b/parallel/ompss2_openacc/set-ompss-openacc-env.sh new file mode 100644 index 0000000..d3d60b9 --- /dev/null +++ b/parallel/ompss2_openacc/set-ompss-openacc-env.sh @@ -0,0 +1,6 @@ +module load nvidia-hpc-sdk/20.11 +module load CUDA/10.2 +module load GCC/10.1.0 +module load BOOST + +export PATH=$PATH:/home/bsc85/bsc85186/ompss_affinity/install/bin/ \ No newline at end of file diff --git a/parallel/ompss2_openacc/simulation.c b/parallel/ompss2_openacc/simulation.c index 1894d07..54fb2b1 100644 --- a/parallel/ompss2_openacc/simulation.c +++ b/parallel/ompss2_openacc/simulation.c @@ -414,7 +414,14 @@ void sim_timings(t_simulation *sim, uint64_t t0, uint64_t t1, const unsigned int fprintf(stdout, "Number of regions (GPU): %d (effective: %d regions)\n", gpu_regions, get_gpu_regions_effective()); fprintf(stdout, "Number of threads: %d\n", n_threads); - fprintf(stdout, "Sort - Bin size: %d\n", TILE_SIZE); + fprintf(stdout, "Number of GPUs: %\n", acc_get_num_devices(acc_device_nvidia)); + +#ifdef ENABLE_AFFINITY + fprintf(stdout, "Affinity Enabled\n"); +#else + fprintf(stdout, "Affinity Disabled\n"); +#endif + // fprintf(stdout, "Time for spec. advance = %f s\n", spec_time() / n_threads); // fprintf(stdout, "Time for emf advance = %f s\n", emf_time() / n_threads); fprintf(stdout, "Total simulation time = %f s\n", timer_interval_seconds(t0, t1));