aboutsummaryrefslogtreecommitdiff
path: root/Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp
blob: e5514c69725b5df55eab34f1f3c10b37742f3347 (plain) (blame)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
// TODO include statements

FiniteDifferenceSolver::EvolveB ( VectorField Bfield,
                                  ConstVectorField Efield,
                                  amrex::Real dt ) {
    // Select algorithm (The choice of algorithm is a runtime option,
    // but we compile code for each algorithm, using templates)
    if (fdtd_algo == MaxwellSolverAlgo::Yee){
        EvolveB <YeeAlgorithm> ( Bfield, Efield, dt );
    } else if (fdtd_algo == MaxwellSolverAlgo::CKC) {
        EvolveB <CKCAlgorithm> ( Bfield, Efield, dt );
    } else {
        amrex::Abort("Unknown algorithm");
    }
)

template<typename algo>
FiniteDifferenceSolver::EvolveB ( VectorField Bfield,
                                 ConstVectorField Efield,
                                 amrex::Real dt ) {

    // Loop through the grids, and over the tiles within each grid
#ifdef _OPENMP
#pragma omp parallel if (Gpu::notInLaunchRegion())
#endif
    for ( MFIter mfi(*Bx, TilingIfNotGPU()); mfi.isValid(); ++mfi ) {

        // Extract field data for this grid/tile
        auto const& Bx = Bfield[0]->array(mfi);
        auto const& By = Bfield[1]->array(mfi);
        auto const& Bz = Bfield[2]->array(mfi);
        auto const& Ex = Efield[0]->array(mfi);
        auto const& Ey = Efield[1]->array(mfi);
        auto const& Ez = Efield[2]->array(mfi);

        // Extract stencil coefficients
        Real const* AMREX_RESTRICT coefs_x = stencil_coefs_x.dataPtr();
        Real const* AMREX_RESTRICT coefs_y = stencil_coefs_y.dataPtr();
        Real const* AMREX_RESTRICT coefs_z = stencil_coefs_z.dataPtr();

        // Extract tileboxes for which to loop
        const Box& tbx  = mfi.tilebox(Bx_nodal_flag);
        const Box& tby  = mfi.tilebox(By_nodal_flag);
        const Box& tbz  = mfi.tilebox(Bz_nodal_flag);

        // Loop over the cells and update the fields
        amrex::ParallelFor(tbx, tby, tbz,

            [=] AMREX_GPU_DEVICE (int i, int j, int k){
                Bx(i, j, k) += dt * algo::UpwardDz( Ey, i, j, k, coefs_z)
                             - dt * algo::UpwardDy( Ez, i, j, k, coefs_y);
            },

            [=] AMREX_GPU_DEVICE (int i, int j, int k){
                By(i, j, k) += dt * algo::UpwardDx( Ez, i, j, k, coefs_x)
                             - dt * algo::UpwardDz( Ex, i, j, k, coefs_z);
            },

            [=] AMREX_GPU_DEVICE (int i, int j, int k){
                Bz(i, j, k) += dt * algo::UpwardDy( Ex, i, j, k, coefs_y)
                             - dt * algo::UpwardDx( Ey, i, j, k, coefs_x);
            }

        );

    }

};