Skip to content

Commit 97a1711

Browse files
authored
Merge pull request #124 from ECP-WarpX/clean_godfrey
Move NCI Godfrey filter to C++
2 parents 83fef5b + ada5eb5 commit 97a1711

File tree

20 files changed

+798
-303
lines changed

20 files changed

+798
-303
lines changed

Examples/Physics_applications/laser_acceleration/inputs.2d.boost

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,8 @@
11
#################################
22
######### BOX PARAMETERS ########
33
#################################
4-
# max_step = 2700
5-
stop_time = 1.9e-12
4+
max_step = 1000
5+
# stop_time = 1.9e-12
66
amr.n_cell = 128 1024
77
amr.max_grid_size = 64
88
amr.blocking_factor = 32
@@ -39,7 +39,7 @@ warpx.serialize_ics = 1
3939
#################################
4040
####### BOOST PARAMETERS ########
4141
#################################
42-
warpx.gamma_boost = 10.
42+
warpx.gamma_boost = 30.
4343
warpx.boost_direction = z
4444
warpx.do_boosted_frame_diagnostic = 1
4545
warpx.num_snapshots_lab = 7
@@ -62,11 +62,11 @@ electrons.momentum_distribution_type = "gaussian"
6262
electrons.xmin = -120.e-6
6363
electrons.xmax = 120.e-6
6464
electrons.zmin = 0.5e-3
65-
electrons.zmax = .0035
65+
electrons.zmax = 1.
6666
electrons.profile = "predefined"
6767
electrons.predefined_profile_name = "parabolic_channel"
6868
# predefined_profile_params = z_start ramp_up plateau ramp_down rc n0
69-
electrons.predefined_profile_params = .5e-3 .5e-3 2.e-3 .5e-3 50.e-6 3.5e24
69+
electrons.predefined_profile_params = .5e-3 .5e-3 2.e-3 .5e-3 50.e-6 3.5e25
7070
electrons.do_continuous_injection = 1
7171

7272
ions.charge = q_e
@@ -77,11 +77,11 @@ ions.momentum_distribution_type = "gaussian"
7777
ions.xmin = -120.e-6
7878
ions.xmax = 120.e-6
7979
ions.zmin = 0.5e-3
80-
ions.zmax = .0035
80+
ions.zmax = 1.
8181
ions.profile = "predefined"
8282
ions.predefined_profile_name = "parabolic_channel"
8383
# predefined_profile_params = z_start ramp_up plateau ramp_down rc n0
84-
ions.predefined_profile_params = .5e-3 .5e-3 2.e-3 .5e-3 50.e-6 3.5e24
84+
ions.predefined_profile_params = .5e-3 .5e-3 2.e-3 .5e-3 50.e-6 3.5e25
8585
ions.do_continuous_injection = 1
8686

8787
beam.charge = -q_e
Lines changed: 38 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,38 @@
1+
#! /usr/bin/env python
2+
3+
import sys
4+
import yt ; yt.funcs.mylog.setLevel(0)
5+
import numpy as np
6+
from scipy import signal
7+
8+
# Build Jx without filter (from other simulation)
9+
my_F_nofilter = np.zeros([16,16])
10+
my_F_nofilter[8,8] = -1.601068065642412e-11
11+
my_F_nofilter[8,7] = -1.601068065642412e-11
12+
13+
# Build 2D filter
14+
filter0 = np.array([.25,.5,.25])
15+
my_order = [1,5]
16+
my_filterx = filter0
17+
my_filtery = filter0
18+
while my_order[0]>1:
19+
my_filterx = np.convolve(my_filterx,filter0)
20+
my_order[0] -= 1
21+
while my_order[1]>1:
22+
my_filtery = np.convolve(my_filtery,filter0)
23+
my_order[1] -= 1
24+
my_filter = my_filterx[:,None]*my_filtery
25+
26+
# Apply filter. my_F_filetered is the theoretical value for filtered field
27+
my_F_filtered = signal.convolve2d(my_F_nofilter, my_filter, boundary='symm', mode='same')
28+
29+
# Get simulation result for F_filtered
30+
filename = sys.argv[1]
31+
ds = yt.load( filename )
32+
sl = yt.SlicePlot(ds, 2, 'jx', aspect=1)
33+
all_data_level_0 = ds.covering_grid(level=0,left_edge=ds.domain_left_edge, dims=ds.domain_dimensions)
34+
F_filtered = all_data_level_0['boxlib', 'jx'].v.squeeze()
35+
36+
# Compare theory and PIC for filtered value
37+
error = np.sum( np.abs(F_filtered - my_F_filtered) ) / np.sum( np.abs(my_F_filtered) )
38+
assert( error < 1.e-14 )
Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,25 @@
1+
max_step = 1
2+
amr.n_cell = 16 16
3+
amr.max_level = 0
4+
amr.blocking_factor = 8
5+
amr.max_grid_size = 8
6+
amr.plot_int = 1
7+
geometry.coord_sys = 0
8+
geometry.is_periodic = 0 0
9+
geometry.prob_lo = -8 -12
10+
geometry.prob_hi = 8 12
11+
warpx.do_pml = 0
12+
algo.current_deposition = 1
13+
algo.charge_deposition = 1
14+
algo.field_gathering = 1
15+
algo.particle_pusher = 0
16+
warpx.cfl = 1.0
17+
18+
particles.nspecies = 1
19+
particles.species_names = electron
20+
electron.charge = -q_e
21+
electron.mass = m_e
22+
electron.injection_style = "SingleParticle"
23+
electron.single_particle_pos = 0.0 0.0 0.0
24+
electron.single_particle_vel = 1.e20 0.0 0.0 # gamma*beta
25+
electron.single_particle_weight = 1.0

Regression/WarpX-tests.ini

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -102,6 +102,21 @@ compileTest = 0
102102
doVis = 0
103103
analysisRoutine = Examples/Modules/nci_corrector/ncicorr_analysis.py
104104

105+
[bilinear_filter]
106+
buildDir = .
107+
inputFile = Examples/Tests/SingleParticle/inputs
108+
runtime_params = warpx.use_filter=1 warpx.filter_npass_each_dir=1 5
109+
dim = 2
110+
addToCompileString =
111+
restartTest = 0
112+
useMPI = 1
113+
numprocs = 2
114+
useOMP = 1
115+
numthreads = 2
116+
compileTest = 0
117+
doVis = 0
118+
analysisRoutine = Examples/Tests/SingleParticle/bilinear_filter_analysis.py
119+
105120
[Langmuir_2d]
106121
buildDir = .
107122
inputFile = Examples/Tests/Langmuir/inputs.rt

Source/Filter/BilinearFilter.H

Lines changed: 7 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -1,30 +1,17 @@
1-
#include <AMReX_MultiFab.H>
2-
#include <AMReX_CudaContainers.H>
1+
#include <Filter.H>
32

4-
#ifndef WARPX_FILTER_H_
5-
#define WARPX_FILTER_H_
3+
#ifndef WARPX_BILIN_FILTER_H_
4+
#define WARPX_BILIN_FILTER_H_
65

7-
class BilinearFilter
6+
class BilinearFilter : public Filter
87
{
98
public:
10-
BilinearFilter () = default;
9+
10+
BilinearFilter() = default;
1111

1212
void ComputeStencils();
13-
void ApplyStencil(amrex::MultiFab& dstmf, const amrex::MultiFab& srcmf, int scomp=0, int dcomp=0, int ncomp=10000);
1413

1514
amrex::IntVect npass_each_dir;
16-
amrex::IntVect stencil_length_each_dir;
17-
18-
// public for cuda
19-
void Filter(const amrex::Box& tbx,
20-
amrex::Array4<amrex::Real const> const& tmp,
21-
amrex::Array4<amrex::Real > const& dst,
22-
int scomp, int dcomp, int ncomp);
23-
24-
private:
25-
26-
amrex::Gpu::ManagedVector<amrex::Real> stencil_x, stencil_y, stencil_z;
2715

28-
amrex::Dim3 slen;
2916
};
30-
#endif
17+
#endif // #ifndef WARPX_BILIN_FILTER_H_

Source/Filter/BilinearFilter.cpp

Lines changed: 0 additions & 170 deletions
Original file line numberDiff line numberDiff line change
@@ -68,173 +68,3 @@ void BilinearFilter::ComputeStencils(){
6868
slen.z = 1;
6969
#endif
7070
}
71-
72-
73-
#ifdef AMREX_USE_CUDA
74-
75-
void
76-
BilinearFilter::ApplyStencil (MultiFab& dstmf, const MultiFab& srcmf, int scomp, int dcomp, int ncomp)
77-
{
78-
BL_PROFILE("BilinearFilter::ApplyStencil()");
79-
ncomp = std::min(ncomp, srcmf.nComp());
80-
81-
for (MFIter mfi(dstmf); mfi.isValid(); ++mfi)
82-
{
83-
const auto& src = srcmf.array(mfi);
84-
const auto& dst = dstmf.array(mfi);
85-
const Box& tbx = mfi.growntilebox();
86-
const Box& gbx = amrex::grow(tbx,stencil_length_each_dir-1);
87-
88-
// tmpfab has enough ghost cells for the stencil
89-
FArrayBox tmp_fab(gbx,ncomp);
90-
Elixir tmp_eli = tmp_fab.elixir(); // Prevent the tmp data from being deleted too early
91-
auto const& tmp = tmp_fab.array();
92-
93-
// Copy values in srcfab into tmpfab
94-
const Box& ibx = gbx & srcmf[mfi].box();
95-
AMREX_PARALLEL_FOR_4D ( gbx, ncomp, i, j, k, n,
96-
{
97-
if (ibx.contains(IntVect(AMREX_D_DECL(i,j,k)))) {
98-
tmp(i,j,k,n) = src(i,j,k,n+scomp);
99-
} else {
100-
tmp(i,j,k,n) = 0.0;
101-
}
102-
});
103-
104-
// Apply filter
105-
Filter(tbx, tmp, dst, 0, dcomp, ncomp);
106-
}
107-
}
108-
109-
void BilinearFilter::Filter (const Box& tbx,
110-
Array4<Real const> const& tmp,
111-
Array4<Real > const& dst,
112-
int scomp, int dcomp, int ncomp)
113-
{
114-
amrex::Real const* AMREX_RESTRICT sx = stencil_x.data();
115-
amrex::Real const* AMREX_RESTRICT sy = stencil_y.data();
116-
amrex::Real const* AMREX_RESTRICT sz = stencil_z.data();
117-
Dim3 slen_local = slen;
118-
AMREX_PARALLEL_FOR_4D ( tbx, ncomp, i, j, k, n,
119-
{
120-
Real d = 0.0;
121-
122-
for (int iz=0; iz < slen_local.z; ++iz){
123-
for (int iy=0; iy < slen_local.y; ++iy){
124-
for (int ix=0; ix < slen_local.x; ++ix){
125-
#if (AMREX_SPACEDIM == 3)
126-
Real sss = sx[ix]*sy[iy]*sz[iz];
127-
#else
128-
Real sss = sx[ix]*sz[iy];
129-
#endif
130-
#if (AMREX_SPACEDIM == 3)
131-
d += sss*( tmp(i-ix,j-iy,k-iz,scomp+n)
132-
+tmp(i+ix,j-iy,k-iz,scomp+n)
133-
+tmp(i-ix,j+iy,k-iz,scomp+n)
134-
+tmp(i+ix,j+iy,k-iz,scomp+n)
135-
+tmp(i-ix,j-iy,k+iz,scomp+n)
136-
+tmp(i+ix,j-iy,k+iz,scomp+n)
137-
+tmp(i-ix,j+iy,k+iz,scomp+n)
138-
+tmp(i+ix,j+iy,k+iz,scomp+n));
139-
#else
140-
d += sss*( tmp(i-ix,j-iy,k,scomp+n)
141-
+tmp(i+ix,j-iy,k,scomp+n)
142-
+tmp(i-ix,j+iy,k,scomp+n)
143-
+tmp(i+ix,j+iy,k,scomp+n));
144-
#endif
145-
}
146-
}
147-
}
148-
149-
dst(i,j,k,dcomp+n) = d;
150-
});
151-
}
152-
153-
#else
154-
155-
void
156-
BilinearFilter::ApplyStencil (MultiFab& dstmf, const MultiFab& srcmf, int scomp, int dcomp, int ncomp)
157-
{
158-
BL_PROFILE("BilinearFilter::ApplyStencil()");
159-
ncomp = std::min(ncomp, srcmf.nComp());
160-
#ifdef _OPENMP
161-
#pragma omp parallel
162-
#endif
163-
{
164-
FArrayBox tmpfab;
165-
for (MFIter mfi(dstmf,true); mfi.isValid(); ++mfi){
166-
const auto& srcfab = srcmf[mfi];
167-
auto& dstfab = dstmf[mfi];
168-
const Box& tbx = mfi.growntilebox();
169-
const Box& gbx = amrex::grow(tbx,stencil_length_each_dir-1);
170-
// tmpfab has enough ghost cells for the stencil
171-
tmpfab.resize(gbx,ncomp);
172-
tmpfab.setVal(0.0, gbx, 0, ncomp);
173-
// Copy values in srcfab into tmpfab
174-
const Box& ibx = gbx & srcfab.box();
175-
tmpfab.copy(srcfab, ibx, scomp, ibx, 0, ncomp);
176-
// Apply filter
177-
Filter(tbx, tmpfab.array(), dstfab.array(), 0, dcomp, ncomp);
178-
}
179-
}
180-
}
181-
182-
void BilinearFilter::Filter (const Box& tbx,
183-
Array4<Real const> const& tmp,
184-
Array4<Real > const& dst,
185-
int scomp, int dcomp, int ncomp)
186-
{
187-
const auto lo = amrex::lbound(tbx);
188-
const auto hi = amrex::ubound(tbx);
189-
// tmp and dst are of type Array4 (Fortran ordering)
190-
amrex::Real const* AMREX_RESTRICT sx = stencil_x.data();
191-
amrex::Real const* AMREX_RESTRICT sy = stencil_y.data();
192-
amrex::Real const* AMREX_RESTRICT sz = stencil_z.data();
193-
for (int n = 0; n < ncomp; ++n) {
194-
// Set dst value to 0.
195-
for (int k = lo.z; k <= hi.z; ++k) {
196-
for (int j = lo.y; j <= hi.y; ++j) {
197-
for (int i = lo.x; i <= hi.x; ++i) {
198-
dst(i,j,k,dcomp+n) = 0.0;
199-
}
200-
}
201-
}
202-
// 3 nested loop on 3D stencil
203-
for (int iz=0; iz < slen.z; ++iz){
204-
for (int iy=0; iy < slen.y; ++iy){
205-
for (int ix=0; ix < slen.x; ++ix){
206-
#if (AMREX_SPACEDIM == 3)
207-
Real sss = sx[ix]*sy[iy]*sz[iz];
208-
#else
209-
Real sss = sx[ix]*sz[iy];
210-
#endif
211-
// 3 nested loop on 3D array
212-
for (int k = lo.z; k <= hi.z; ++k) {
213-
for (int j = lo.y; j <= hi.y; ++j) {
214-
AMREX_PRAGMA_SIMD
215-
for (int i = lo.x; i <= hi.x; ++i) {
216-
#if (AMREX_SPACEDIM == 3)
217-
dst(i,j,k,dcomp+n) += sss*(tmp(i-ix,j-iy,k-iz,scomp+n)
218-
+tmp(i+ix,j-iy,k-iz,scomp+n)
219-
+tmp(i-ix,j+iy,k-iz,scomp+n)
220-
+tmp(i+ix,j+iy,k-iz,scomp+n)
221-
+tmp(i-ix,j-iy,k+iz,scomp+n)
222-
+tmp(i+ix,j-iy,k+iz,scomp+n)
223-
+tmp(i-ix,j+iy,k+iz,scomp+n)
224-
+tmp(i+ix,j+iy,k+iz,scomp+n));
225-
#else
226-
dst(i,j,k,dcomp+n) += sss*(tmp(i-ix,j-iy,k,scomp+n)
227-
+tmp(i+ix,j-iy,k,scomp+n)
228-
+tmp(i-ix,j+iy,k,scomp+n)
229-
+tmp(i+ix,j+iy,k,scomp+n));
230-
#endif
231-
}
232-
}
233-
}
234-
}
235-
}
236-
}
237-
}
238-
}
239-
240-
#endif

Source/Filter/Filter.H

Lines changed: 43 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,43 @@
1+
#include <AMReX_MultiFab.H>
2+
#include <AMReX_CudaContainers.H>
3+
4+
#ifndef WARPX_FILTER_H_
5+
#define WARPX_FILTER_H_
6+
7+
class Filter
8+
{
9+
public:
10+
Filter () = default;
11+
12+
// Apply stencil on MultiFab.
13+
// Guard cells are handled inside this function
14+
void ApplyStencil(amrex::MultiFab& dstmf,
15+
const amrex::MultiFab& srcmf, int scomp=0,
16+
int dcomp=0, int ncomp=10000);
17+
18+
// Apply stencil on a FabArray.
19+
void ApplyStencil (amrex::FArrayBox& dstfab,
20+
const amrex::FArrayBox& srcfab, const amrex::Box& tbx,
21+
int scomp=0, int dcomp=0, int ncomp=10000);
22+
23+
// public for cuda
24+
void DoFilter(const amrex::Box& tbx,
25+
amrex::Array4<amrex::Real const> const& tmp,
26+
amrex::Array4<amrex::Real > const& dst,
27+
int scomp, int dcomp, int ncomp);
28+
29+
// In 2D, stencil_length_each_dir = {length(stencil_x), length(stencil_z)}
30+
amrex::IntVect stencil_length_each_dir;
31+
32+
protected:
33+
// Stencil along each direction.
34+
// in 2D, stencil_y is not initialized.
35+
amrex::Gpu::ManagedVector<amrex::Real> stencil_x, stencil_y, stencil_z;
36+
// Length of each stencil.
37+
// In 2D, slen = {length(stencil_x), length(stencil_z), 1}
38+
amrex::Dim3 slen;
39+
40+
private:
41+
42+
};
43+
#endif // #ifndef WARPX_FILTER_H_

0 commit comments

Comments
 (0)