WarpX
Loading...
Searching...
No Matches
MatrixPC.H
Go to the documentation of this file.
1/* Copyright 2024 Debojyoti Ghosh
2 *
3 * This file is part of WarpX.
4 *
5 * License: BSD-3-Clause-LBNL
6 */
7#ifndef MATRIX_PC_H_
8#define MATRIX_PC_H_
9
10#include "Fields.H"
11#include "Utils/WarpXConst.H"
12#include "Utils/TextMsg.H"
13#include "Preconditioner.H"
14
16
17#include <AMReX.H>
18#include <AMReX_ParmParse.H>
19#include <AMReX_Array.H>
20#include <AMReX_Vector.H>
21#include <AMReX_MultiFab.H>
22
24{
26 bool insertOrAdd( const int a_cidx,
27 const amrex::Real a_val,
28 int* const a_cidxs,
29 amrex::Real* const a_aij,
30 const int a_nnz,
31 int& a_ncol )
32 {
33 if (a_cidx < 0) { return true; /* outside domain */ }
34 int loc = -1;
35 for (int icol = 0; icol < std::min(a_ncol,a_nnz); icol++) {
36 if (a_cidxs[icol] == a_cidx) {
37 loc = icol;
38 break;
39 }
40 }
41 if (loc < 0) {
42 a_ncol++;
43 if (a_ncol > a_nnz) { return false; }
44 else {
45 // column index not found; add new entry
46 a_cidxs[a_ncol-1] = a_cidx;
47 a_aij[a_ncol-1] = a_val;
48 }
49 } else {
50 // column index already exists; add to value
51 a_aij[loc] += a_val;
52 }
53 return true;
54 }
55}
56
74
75template <class T, class Ops>
76class MatrixPC : public Preconditioner<T,Ops>
77{
78 public:
79
80 using RT = typename T::value_type;
81
85 MatrixPC () = default;
86
90 ~MatrixPC () override = default;
91
92 // Prohibit move and copy operations
93 MatrixPC(const MatrixPC&) = delete;
94 MatrixPC& operator=(const MatrixPC&) = delete;
95 MatrixPC(MatrixPC&&) noexcept = delete;
96 MatrixPC& operator=(MatrixPC&&) noexcept = delete;
97
101 void Define (const T&, Ops*) override;
102
106 void Update (const T& a_U) override;
107
116 int Assemble (const T& a_U);
117
126 void Apply (T&, const T&) override;
127
128 inline void getPCMatrix (amrex::Gpu::DeviceVector<int>& a_r_indices_g,
129 amrex::Gpu::DeviceVector<int>& a_num_nz,
130 amrex::Gpu::DeviceVector<int>& a_c_indices_g,
131 amrex::Gpu::DeviceVector<RT>& a_a_ij,
132 int& a_n, int& a_ncols_max ) override
133 {
134 a_n = m_ndofs_l;
135 a_ncols_max = m_pc_mat_nnz;
136
137 a_r_indices_g.resize(m_r_indices_g.size());
139 m_r_indices_g.begin(),
140 m_r_indices_g.end(),
141 a_r_indices_g.begin() );
142
143 a_num_nz.resize(m_num_nz.size());
145 m_num_nz.begin(),
146 m_num_nz.end(),
147 a_num_nz.begin() );
148
149 a_c_indices_g.resize(m_c_indices_g.size());
151 m_c_indices_g.begin(),
152 m_c_indices_g.end(),
153 a_c_indices_g.begin() );
154
155 a_a_ij.resize(m_a_ij.size());
157 m_a_ij.begin(),
158 m_a_ij.end(),
159 a_a_ij.begin() );
161 }
162
166 void printParameters() const override;
167
171 [[nodiscard]] inline bool IsDefined () const override { return m_is_defined; }
172
173 inline void setName (const std::string& a_name) override { m_name = a_name; }
174
175 protected:
176
177 bool m_is_defined = false;
178 bool m_verbose = true;
179
180 Ops* m_ops = nullptr;
181
184
185 int m_ndofs_l = 0;
186 int m_ndofs_g = 0;
187 bool m_pc_diag_only = false;
190
191 std::string m_name = "noname";
192
197
199
201
205 void readParameters();
206
207 private:
208
209};
210
211template <class T, class Ops>
213{
214 using namespace amrex;
215 Print() << m_name << " verbose: " << (m_verbose?"true":"false") << "\n";
216 Print() << m_name << " pc_diagonal_only: " << (m_pc_diag_only?"true":"false") << "\n";
217 Print() << m_name << " include_mass_matrices: " << (m_include_mass_matrices?"true":"false") << "\n";
218}
219
220template <class T, class Ops>
222{
224 pp.query("verbose", m_verbose);
225 pp.query("pc_diagonal_only", m_pc_diag_only);
226}
227
228template <class T, class Ops>
229void MatrixPC<T,Ops>::Define ( const T& a_U,
230 Ops* const a_ops )
231{
232 BL_PROFILE("MatrixPC::Define()");
233 using namespace amrex;
234
236 !IsDefined(),
237 "MatrixPC::Define() called on defined object" );
239 (a_ops != nullptr),
240 "MatrixPC::Define(): a_ops is nullptr" );
242 a_U.getArrayVecType()==warpx::fields::FieldType::Efield_fp,
243 "MatrixPC::Define() must be called with Efield_fp type");
244
245 m_ops = a_ops;
246 // read preconditioner parameters
248
249 // a_U is not needed
251
252 // Set number of AMR levels and create geometry, grids, and
253 // distribution mapping vectors.
254 m_num_amr_levels = m_ops->numAMRLevels();
255 if (m_num_amr_levels > 1) {
256 WARPX_ABORT_WITH_MESSAGE("MatrixPC::Define(): m_num_amr_levels > 1");
257 }
258 m_geom.resize(m_num_amr_levels);
259 for (int n = 0; n < m_num_amr_levels; n++) {
260 m_geom[n] = m_ops->GetGeometry(n);
261 }
262
263 m_ndofs_l = a_U.nDOF_local();
264 m_ndofs_g = a_U.nDOF_global();
265
266 auto n_rows = size_t(m_ndofs_l);
267 auto n_cols = size_t(m_pc_mat_nnz) * size_t(m_ndofs_l);
268
269 m_r_indices_g.resize(n_rows);
270 m_num_nz.resize(n_rows);
271 m_c_indices_g.resize(n_cols);
272 m_a_ij.resize(n_cols);
273
274
275 m_bcoefs = m_ops->GetMassMatricesCoeff();
276 if (m_bcoefs != nullptr) {
278 }
279
280 m_is_defined = true;
281}
282
283template <class T, class Ops>
284void MatrixPC<T,Ops>::Update (const T& a_U)
285{
286 BL_PROFILE("MatrixPC::Update()");
287 using namespace amrex;
288
290 IsDefined(),
291 "MatrixPC::Update() called on undefined object" );
292
293 while(true) {
294
295 auto nnz_diff = Assemble(a_U);
296 AMREX_ALWAYS_ASSERT(nnz_diff >= 0);
297 if (nnz_diff) {
298
299 m_pc_mat_nnz += nnz_diff;
301
302 } else {
303 break;
304 }
305 }
306
307 if (m_num_realloc > 1) {
308 std::stringstream warning_message;
309 warning_message << "Number of times arrays were reallocated due to new nonzero elements "
310 << "is greater than 1 (" << m_num_realloc <<"). This is unexpected.\n";
311 ablastr::warn_manager::WMRecordWarning("MatrixPC", warning_message.str());
312 }
313
314}
315
316template <class T, class Ops>
318{
319 // Assemble the sparse matrix representation of the preconditioner
320 // A = curl (alpha * curl []) + M
321 // where M is the mass matrix. The following data is set in this function:
322 // - m_r_indices_g: integer array of size n with the global row indices
323 // - m_num_nz: integer array of size n with the number of non-zero elements
324 // in each row
325 // - m_c_indices: integer array of size n*ncmax with the global column indices
326 // of non-zero elements in each row (row-major)
327 // - m_a_ij: real-type array of size n*ncmax with the matrix element values
328 // (row-major format)
329 // where n is the local number of rows, and ncmax is the maximum number of non-zero
330 // elements per row.
331
332 BL_PROFILE("MatrixPC::Assemble()");
333 using namespace amrex;
334
336 IsDefined(),
337 "MatrixPC::Assemble() called on undefined object" );
338
339 // set the alpha coefficient for the curl-curl op
340 const RT thetaDt = m_ops->GetThetaForPC()*this->m_dt;
341 const RT alpha = (thetaDt*PhysConst::c) * (thetaDt*PhysConst::c);
342 if (m_verbose) {
343 amrex::Print() << "Updating " << m_name
344 << ": theta*dt = " << thetaDt << ", "
345 << " coefficients: "
346 << "alpha = " << alpha << "\n";
347 }
348
349 // Get DOF object from a_U
350 const auto& dofs_obj = a_U.getDOFsObject();
351 const auto& dofs_mfarrvec = dofs_obj->m_array;
352 AMREX_ALWAYS_ASSERT(m_ndofs_l == dofs_obj->m_nDoFs_l);
353 AMREX_ALWAYS_ASSERT(m_ndofs_g == dofs_obj->m_nDoFs_g);
354
355 m_r_indices_g.clear();
356 m_num_nz.clear();
357 m_c_indices_g.clear();
358 m_a_ij.clear();
359
360 auto n_rows = size_t(m_ndofs_l);
361 auto n_cols = size_t(m_pc_mat_nnz) * size_t(m_ndofs_l);
362
363 m_r_indices_g.resize(n_rows);
364 m_num_nz.resize(n_rows);
365 m_c_indices_g.resize(n_cols);
366 m_a_ij.resize(n_cols);
367
368 auto* r_indices_g_ptr = m_r_indices_g.data();
369 auto* num_nz_ptr = m_num_nz.data();
370 auto* c_indices_g_ptr = m_c_indices_g.data();
371 auto* a_ij_ptr = m_a_ij.data();
372
373 const auto nnz_max = m_pc_mat_nnz;
374 auto nnz_actual = nnz_max;
375
376 for (int lev = 0; lev < m_num_amr_levels; lev++) {
377
378 auto ncomp = dofs_mfarrvec[lev][0]->nComp();
379 AMREX_ALWAYS_ASSERT(ncomp == 2); // local, global
380
381 const auto& geom = m_geom[lev];
382 const auto dxi = geom.InvCellSizeArray();
383
384 Gpu::Buffer<int> nnz_actual_d({nnz_max});
385 auto* nnz_actual_ptr = nnz_actual_d.data();
386
387 for (int dir = 0; dir < 3; dir++) {
388
389 for (amrex::MFIter mfi(*dofs_mfarrvec[lev][dir]); mfi.isValid(); ++mfi) {
390
391 const amrex::Box bx = mfi.tilebox();
392 const amrex::Box full_bx = mfi.fabbox();
393
394 auto dof_arr = dofs_mfarrvec[lev][dir]->const_array(mfi);
395
396 // Set row indices and identity diagonal (unconditional)
397 ParallelFor(bx, [=] AMREX_GPU_DEVICE (int i, int j, int k)
398 {
399 const int ridx_l = dof_arr(i,j,k,0);
400 if (ridx_l < 0) { return; }
401
402 int icol = 0;
403 const int ridx_g = dof_arr(i,j,k,1);
404
405 r_indices_g_ptr[ridx_l] = ridx_g;
406
407 {
408 const int cidx_g_lhs = dof_arr(i,j,k,1);
409 const amrex::Real val = 1.0_rt;
410 auto flag = MatrixPCUtils::insertOrAdd( cidx_g_lhs, val,
411 &c_indices_g_ptr[ridx_l*nnz_max],
412 &a_ij_ptr[ridx_l*nnz_max],
413 nnz_max, icol );
414 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
415 }
416
417 num_nz_ptr[ridx_l] = icol;
418 });
419
420 // Add the curl-curl stencil entries (only when alpha > 0)
421 if (thetaDt > 0.0) {
422
423#if defined(WARPX_DIM_RSPHERE)
424 // 1D spherical geometry is electrostatic
426#else
427 const amrex::MultiFab* BC_mask_Edir = m_ops->GetCurl2BCmask(lev,dir);
428 AMREX_ALWAYS_ASSERT(BC_mask_Edir != nullptr);
429 const auto BC_mask_Edir_arr = BC_mask_Edir->const_array(mfi);
430#endif
431
432#if defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ)
433 int tdir = -1;
434 if (dir == 0) { tdir = 2; }
435 else if (dir == 2) { tdir = 0; }
436 else { tdir = 1; }
437 auto dof_tdir_arr = dofs_mfarrvec[lev][tdir]->const_array(mfi);
438#elif defined(WARPX_DIM_3D)
439 const int tdir1 = (dir + 1) % 3;
440 const int tdir2 = (dir + 2) % 3;
441 GpuArray<Array4<const int>, AMREX_SPACEDIM>
442 const dof_arrays {{ AMREX_D_DECL( dofs_mfarrvec[lev][dir]->const_array(mfi),
443 dofs_mfarrvec[lev][tdir1]->const_array(mfi),
444 dofs_mfarrvec[lev][tdir2]->const_array(mfi) ) }};
445#endif
446
447 ParallelFor(bx, [=] AMREX_GPU_DEVICE (int i, int j, int k)
448 {
449 const int ridx_l = dof_arr(i,j,k,0);
450 if (ridx_l < 0) { return; }
451
452 int icol = num_nz_ptr[ridx_l]; //NOLINT (misc-const-correctness)
453
454#if defined(WARPX_DIM_1D_Z)
455 // dir = 0: xhat\cdot[\nabla\times\nabla E] = -[d2/dx2]Ex
456 // dir = 1: yhat\cdot[\nabla\times\nabla E] = -[d2/dx2]Ey
457 // dir = 2: zhat\cdot[\nabla\times\nabla E] = 0
458 if (dir != 2) {
459 {
460 // diagonal component (i) of 2nd derivative operator
461 const int cidx_g_rhs = dof_arr(i,j,k,1);
462 const amrex::Real val = 2.0_rt*alpha * dxi[0]*dxi[0] * BC_mask_Edir_arr(i,j,k,0);
463 const auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
464 &c_indices_g_ptr[ridx_l*nnz_max],
465 &a_ij_ptr[ridx_l*nnz_max],
466 nnz_max, icol );
467 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
468 }
469 {
470 // left component (i-1) of 2nd derivative operator
471 const int cidx_g_rhs = dof_arr(i-1,j,k,1);
472 const amrex::Real val = -alpha * dxi[0]*dxi[0] * BC_mask_Edir_arr(i,j,k,1);
473 const auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
474 &c_indices_g_ptr[ridx_l*nnz_max],
475 &a_ij_ptr[ridx_l*nnz_max],
476 nnz_max, icol );
477 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
478 }
479 {
480 // right component (i+1) of 2nd derivative operator
481 const int cidx_g_rhs = dof_arr(i+1,j,k,1);
482 const amrex::Real val = -alpha * dxi[0]*dxi[0] * BC_mask_Edir_arr(i,j,k,1);
483 const auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
484 &c_indices_g_ptr[ridx_l*nnz_max],
485 &a_ij_ptr[ridx_l*nnz_max],
486 nnz_max, icol );
487 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
488 }
489 }
490#elif defined(WARPX_DIM_RCYLINDER)
491 // dir = 0: rhat\cdot[\nabla\times\nabla E] = 0
492 // dir = 1: that\cdot[\nabla\times\nabla E] = -d/dr[1/r*d/dr(r*Et)]
493 // dir = 2: zhat\cdot[\nabla\times\nabla E] = -1/r*d/dr[r*dEz/dr]
494 if (dir != 0) {
495 const auto i_real = static_cast<amrex::Real>(i);
496 {
497 // diagonal component (i) of 2nd derivative operator
498 int cidx_g_rhs = dof_arr(i,j,k,1);
499 amrex::Real geom_factor = 1.0_rt;
500 if (dir == 1) {
501 // r_{i} / r_{i-1/2} + r_{i} / r_{i+1/2}
502 geom_factor = i_real / (i_real - 0.5_rt) + i_real / (i_real + 0.5_rt);
503 }
504 else if (dir == 2) {
505 geom_factor = 2.0_rt; // (r_{i+1/2} + r_{i-1/2}) / r_{i}
506 }
507 amrex::Real val = geom_factor * alpha * dxi[0]*dxi[0] * BC_mask_Edir_arr(i,j,k,0);
508 auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
509 &c_indices_g_ptr[ridx_l*nnz_max],
510 &a_ij_ptr[ridx_l*nnz_max],
511 nnz_max, icol );
512 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
513 }
514 {
515 // left component (i-1) of 2nd derivative operator
516 int cidx_g_rhs = dof_arr(i-1,j,k,1);
517 amrex::Real geom_factor = 1.0_rt;
518 if (dir == 1) {
519 geom_factor = (i_real - 1.0_rt) / (i_real - 0.5_rt); // r_{i-1} / r_{i-1/2}
520 }
521 else if (dir == 2 && i != 0) {
522 geom_factor = 1.0_rt - 0.5_rt / i_real; // r_{i-1/2} / r_{i}
523 }
524 amrex::Real val = -geom_factor * alpha * dxi[0]*dxi[0] * BC_mask_Edir_arr(i,j,k,1);
525 auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
526 &c_indices_g_ptr[ridx_l*nnz_max],
527 &a_ij_ptr[ridx_l*nnz_max],
528 nnz_max, icol );
529 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
530 }
531 {
532 // right component (i+1) of 2nd derivative operator
533 int cidx_g_rhs = dof_arr(i+1,j,k,1);
534 amrex::Real geom_factor = 1.0_rt;
535 if (dir == 1) {
536 geom_factor = (i_real + 1.0_rt) / (i_real + 0.5_rt); // r_{i+1} / r_{i+1/2}
537 }
538 else if (dir == 2 && i != 0) {
539 geom_factor = 1.0_rt + 0.5_rt / i_real; // r_{i+1/2} / r_{i}
540 }
541 amrex::Real val = -geom_factor * alpha * dxi[0]*dxi[0] * BC_mask_Edir_arr(i,j,k,1);
542 auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
543 &c_indices_g_ptr[ridx_l*nnz_max],
544 &a_ij_ptr[ridx_l*nnz_max],
545 nnz_max, icol );
546 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
547 }
548 }
549#elif defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ)
550 // DIM_XZ:
551 // dir = 0: xhat\cdot[\nabla\times\nabla E] = d/dx[dEz/dz] - [d2/dz2]Ex
552 // dir = 1: yhat\cdot[\nabla\times\nabla E] = -[d2/dx2 + d2/dz2]Ey
553 // dir = 2: zhat\cdot[\nabla\times\nabla E] = d/dz[dEx/dx] - [d2/dx2]Ez
554 // DIM_RZ:
555 // dir = 0: rhat\cdot[\nabla\times\nabla E] = d/dr[dEz/dz] - [d2/dz2]Er
556 // dir = 1: that\cdot[\nabla\times\nabla E] = -[d2/dz2]Et - d/dr[1/r*d/dr(r*Et)]
557 // dir = 2: zhat\cdot[\nabla\times\nabla E] = 1/r*d/dr[r*dEr/dz] - 1/r*d/dr[r*dEz/dr]
558#if defined(WARPX_DIM_RZ)
559 const auto i_real = static_cast<amrex::Real>(i);
560#endif
561 {
562 // diagonal component (i,j) of second derivative operator
563 const int cidx_g_rhs = dof_arr(i,j,k,1);
564 amrex::Real val = 0.0_rt;
565 if (dir == 0) {
566 val = 2.0_rt * alpha * dxi[1]*dxi[1] * BC_mask_Edir_arr(i,j,k,0);
567 } else if (dir == 2) {
568 val = 2.0_rt * alpha * dxi[0]*dxi[0] * BC_mask_Edir_arr(i,j,k,0);
569 } else if (dir == 1) {
570#if defined(WARPX_DIM_RZ)
571 const amrex::Real geom_factor = i_real / (i_real - 0.5_rt) + i_real / (i_real + 0.5_rt);
572#else
573 const amrex::Real geom_factor = 2.0_rt;
574#endif
575 val = geom_factor * alpha * dxi[0]*dxi[0] * BC_mask_Edir_arr(i,j,k,0)
576 + 2.0_rt * alpha * dxi[1]*dxi[1] * BC_mask_Edir_arr(i,j,k,2);
577 }
578 auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
579 &c_indices_g_ptr[ridx_l*nnz_max],
580 &a_ij_ptr[ridx_l*nnz_max],
581 nnz_max, icol );
582 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
583 }
584 if ((dir == 0) || (dir == 2)) {
585 {
586 // left component (i-1,j) or (i,j-1) of second derivative operator
587#if defined(WARPX_DIM_RZ)
588 const amrex::Real geom_factor = (dir == 2 && i != 0 ? 1.0_rt - 0.5_rt / i_real : 1.0_rt);
589#else
590 const amrex::Real geom_factor = 1.0_rt;
591#endif
592 const int cidx_g_rhs = (dir == 0 ? dof_arr(i,j-1,k,1) : dof_arr(i-1,j,k,1));
593 const amrex::Real val = -geom_factor * alpha * dxi[dir==0?1:0] * dxi[dir==0?1:0] * BC_mask_Edir_arr(i,j,k,1);
594 const auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
595 &c_indices_g_ptr[ridx_l*nnz_max],
596 &a_ij_ptr[ridx_l*nnz_max],
597 nnz_max, icol );
598 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
599 }
600 {
601 // right component (i+1,j) or (i,j+1) of second derivative operator
602#if defined(WARPX_DIM_RZ)
603 const amrex::Real geom_factor = (dir == 2 && i != 0 ? 1.0_rt + 0.5_rt / i_real : 1.0_rt);
604#else
605 const amrex::Real geom_factor = 1.0_rt;
606#endif
607 const int cidx_g_rhs = (dir == 0 ? dof_arr(i,j+1,k,1) : dof_arr(i+1,j,k,1));
608 const amrex::Real val = -geom_factor * alpha * dxi[dir==0?1:0] * dxi[dir==0?1:0] * BC_mask_Edir_arr(i,j,k,1);
609 const auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
610 &c_indices_g_ptr[ridx_l*nnz_max],
611 &a_ij_ptr[ridx_l*nnz_max],
612 nnz_max, icol );
613 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
614 }
615 // The following four blocks are for the mixed derivative terms of the curl curl operator
616 // DIM_XZ:
617 // dir = 0: d/dx(dEz/dz) at Ex(i,j) with Ex centered in x and nodal in z
618 // dir = 2: d/dz(dEx/dx) at Ez(i,j) with Ez centered in z and nodal in x
619 // DIM_RZ:
620 // dir = 0: d/dr(dEz/dz) at Er(i,j) with Er centered in r and nodal in z
621 // dir = 2: 1/r*d/dr[r*dEr/dz] at Ez(i,j) with Ez centered in z and nodal in r
622#if defined(WARPX_DIM_RZ)
623 const amrex::Real geom_m = (dir == 2 && i != 0 ? 1.0_rt - 0.5_rt / i_real : 1.0_rt);
624 const amrex::Real geom_p = (dir == 2 && i != 0 ? 1.0_rt + 0.5_rt / i_real : 1.0_rt);
625#else
626 const amrex::Real geom_m = 1.0_rt;
627 const amrex::Real geom_p = 1.0_rt;
628#endif
629 {
630 const int cidx_g_rhs = (dir == 0 ? dof_tdir_arr(i,j-1,k,1) : dof_tdir_arr(i-1,j,k,1));
631 const amrex::Real val = geom_m * alpha * dxi[0] * dxi[1] * BC_mask_Edir_arr(i,j,k,2);
632 const auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
633 &c_indices_g_ptr[ridx_l*nnz_max],
634 &a_ij_ptr[ridx_l*nnz_max],
635 nnz_max, icol );
636 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
637 }
638 {
639 const int cidx_g_rhs = dof_tdir_arr(i,j,k,1);
640 const amrex::Real val = -geom_p * alpha * dxi[0] * dxi[1] * BC_mask_Edir_arr(i,j,k,2);
641 const auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
642 &c_indices_g_ptr[ridx_l*nnz_max],
643 &a_ij_ptr[ridx_l*nnz_max],
644 nnz_max, icol );
645 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
646 }
647 {
648 const int cidx_g_rhs = (dir == 0 ? dof_tdir_arr(i+1,j-1,k,1) : dof_tdir_arr(i-1,j+1,k,1));
649 const amrex::Real val = -geom_m * alpha * dxi[0] * dxi[1] * BC_mask_Edir_arr(i,j,k,2);
650 const auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
651 &c_indices_g_ptr[ridx_l*nnz_max],
652 &a_ij_ptr[ridx_l*nnz_max],
653 nnz_max, icol );
654 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
655 }
656 {
657 const int cidx_g_rhs = (dir == 0 ? dof_tdir_arr(i+1,j,k,1) : dof_tdir_arr(i,j+1,k,1));
658 const amrex::Real val = geom_p * alpha * dxi[0] * dxi[1] * BC_mask_Edir_arr(i,j,k,2);
659 const auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
660 &c_indices_g_ptr[ridx_l*nnz_max],
661 &a_ij_ptr[ridx_l*nnz_max],
662 nnz_max, icol );
663 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
664 }
665 } else if (dir==1) {
666 for (int jdir = 0; jdir <= 2; jdir+=2) {
667 {
668#if defined(WARPX_DIM_RZ)
669 const amrex::Real geom_factor = (jdir == 0 ? (i_real - 1.0_rt) / (i_real - 0.5_rt) : 1.0_rt);
670#else
671 const amrex::Real geom_factor = 1.0_rt;
672#endif
673 const int cidx_g_rhs = (jdir == 0 ? dof_arr(i-1,j,k,1) : dof_arr(i,j-1,k,1));
674 const amrex::Real val = -geom_factor * alpha * dxi[jdir==0?0:1] * dxi[jdir==0?0:1] * BC_mask_Edir_arr(i,j,k,jdir+1);
675 const auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
676 &c_indices_g_ptr[ridx_l*nnz_max],
677 &a_ij_ptr[ridx_l*nnz_max],
678 nnz_max, icol );
679 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
680 }
681 {
682#if defined(WARPX_DIM_RZ)
683 const amrex::Real geom_factor = (jdir == 0 ? (i_real + 1.0_rt) / (i_real + 0.5_rt) : 1.0_rt);
684#else
685 const amrex::Real geom_factor = 1.0_rt;
686#endif
687 const int cidx_g_rhs = (jdir == 0 ? dof_arr(i+1,j,k,1) : dof_arr(i,j+1,k,1));
688 const amrex::Real val = -geom_factor * alpha * dxi[jdir==0?0:1] * dxi[jdir==0?0:1] * BC_mask_Edir_arr(i,j,k,jdir+1);
689 const auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
690 &c_indices_g_ptr[ridx_l*nnz_max],
691 &a_ij_ptr[ridx_l*nnz_max],
692 nnz_max, icol );
693 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
694 }
695 }
696 }
697#elif defined(WARPX_DIM_3D)
698 // xhat\cdot[\nabla\times\nabla E] = d/dx[dEy/dy + dEz/dz] - [d2/dy2 + d2/dz2]Ex
699 // yhat\cdot[\nabla\times\nabla E] = d/dy[dEx/dx + dEz/dz] - [d2/dx2 + d2/dz2]Ey
700 // zhat\cdot[\nabla\times\nabla E] = d/dz[dEx/dx + dEy/dy] - [d2/dx2 + d2/dy2]Ez
701 amrex::IntVect dvec(AMREX_D_DECL(dir,tdir1,tdir2));
702 const amrex::IntVect ic(AMREX_D_DECL(i,j,k));
703 {
704 // diagonal component (i,j,k) of second derivative operator
705 const int cidx_g_rhs = dof_arrays[0](ic,1);
706 const amrex::Real val = 2.0_rt * alpha * ( dxi[dvec[1]]*dxi[dvec[1]] * BC_mask_Edir_arr(i,j,k,0)
707 + dxi[dvec[2]]*dxi[dvec[2]] * BC_mask_Edir_arr(i,j,k,3) );
708 const auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
709 &c_indices_g_ptr[ridx_l*nnz_max],
710 &a_ij_ptr[ridx_l*nnz_max],
711 nnz_max, icol );
712 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
713 }
714 // left and right components (i+/-1,j,k), (i,j+/-1,k), or (i,j,k+/-1) of second derivative operator
715 for (int ctr = -1; ctr <= 1; ctr += 2) {
716 for (int tdir = 1; tdir <= 2; tdir++) {
717 auto iv = ic; iv[dvec[tdir]] += ctr;
718 const int cidx_g_rhs = dof_arrays[0](iv,1);
719 const int comp_shift = (dvec[tdir] == tdir1) ? 0:3;
720 const amrex::Real val = -alpha * dxi[dvec[tdir]]*dxi[dvec[tdir]] * BC_mask_Edir_arr(i,j,k,comp_shift+1);
721 const auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
722 &c_indices_g_ptr[ridx_l*nnz_max],
723 &a_ij_ptr[ridx_l*nnz_max],
724 nnz_max, icol );
725 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
726 }
727 }
728 // mixed derivative terms of the curl curl operator
729 for (int ctr_dir = -1; ctr_dir <= 0; ctr_dir++) {
730 for (int ctr_tdir = -1; ctr_tdir <= 0; ctr_tdir++) {
731 for (int tdir = 1; tdir <= 2; tdir++) {
732 auto iv = ic; iv[dvec[0]] += (ctr_dir+1); iv[dvec[tdir]] += ctr_tdir;
733 const auto sign = std::copysign(1,ctr_dir) * std::copysign(1,ctr_tdir);
734 const int cidx_g_rhs = dof_arrays[tdir](iv,1);
735 const int comp_shift = (dvec[tdir] == tdir1) ? 0:3;
736 const amrex::Real val = amrex::Real(sign) * alpha * dxi[dvec[0]]*dxi[dvec[tdir]] * BC_mask_Edir_arr(i,j,k,comp_shift+2);
737 const auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
738 &c_indices_g_ptr[ridx_l*nnz_max],
739 &a_ij_ptr[ridx_l*nnz_max],
740 nnz_max, icol );
741 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
742 }
743 }
744 }
745#endif
746 num_nz_ptr[ridx_l] = icol;
747 });
748 }
749
750 // Add the mass matrix piece
751 // See Figure B.9 of JCP 491, 112383 (2023) for an illustrative diagram
752 // of the mass matrices (https://doi.org/10.1016/j.jcp.2023.112383).
753 //
754 // The coupling of Jx(i,j,k) to Ex(i+i0,j+j0,k+k0), where i0 ranges from
755 // -MM_width[0] to +MM_width[0], j0 ranges from -MM_width[1] to +MM_width[1], and k0 ranges
756 // from -MM_width[2] to +MM_width[2], is stored as components of m_bcoefs[dir=0] (mass matrices).
757 // Similarly for Jy/Ey (m_bcoefs[dir=1]) and Jz/Ez (m_bcoefs[dir=2]).
758 // The mapping to the components is given by:
759 // mm_comp = i0 + MM_width[0] + MM_comp[0]*(j0 + MM_width[1]) + (MM_comp[0] + MM_comp[1])*(k0 + MM_width[2])
761
762 auto sigma_ii_arr = (*m_bcoefs)[lev][dir]->const_array(mfi);
763 amrex::GpuArray<int,3> MM_ncomp = {1,1,1};
764 amrex::GpuArray<int,3> MM_width = {0,0,0};
765 for (int space_dir=0; space_dir<AMREX_SPACEDIM; space_dir++) {
766 MM_ncomp[space_dir] = m_ops->GetMassMatricesPCnComp(dir,space_dir);
767 MM_width[space_dir] = (MM_ncomp[space_dir] - 1)/2;
768 }
769
770 ParallelFor(bx, [=] AMREX_GPU_DEVICE (int i, int j, int k)
771 {
772 const int ridx_l = dof_arr(i,j,k,0);
773 if (ridx_l < 0) { return; }
774
775 const amrex::IntVect iv_base = IntVect(AMREX_D_DECL(i,j,k));
776 int icol = num_nz_ptr[ridx_l];
777
778 int mm_comp = 0;
779 for (int comp2 = 0; comp2 < MM_ncomp[2]; comp2++) {
780 [[maybe_unused]] const int kk0 = comp2 - MM_width[2];
781 for (int comp1 = 0; comp1 < MM_ncomp[1]; comp1++) {
782 [[maybe_unused]] const int jj0 = comp1 - MM_width[1];
783 for (int comp0 = 0; comp0 < MM_ncomp[0]; comp0++) {
784 const int ii0 = comp0 - MM_width[0]; // ..., -2, -1, 0, 1, 2, ...
785 const amrex::IntVect iv_shift = IntVect(AMREX_D_DECL(ii0, jj0, kk0));
786 if (full_bx.contains(iv_base + iv_shift)) {
787 const int cidx_g_rhs = dof_arr(iv_base + iv_shift,1);
788 const amrex::Real val = sigma_ii_arr(iv_base,mm_comp);
789 auto flag = MatrixPCUtils::insertOrAdd( cidx_g_rhs, val,
790 &c_indices_g_ptr[ridx_l*nnz_max],
791 &a_ij_ptr[ridx_l*nnz_max],
792 nnz_max, icol );
793 if (!flag) { Gpu::Atomic::Max(nnz_actual_ptr, icol); }
794 }
795 ++mm_comp;
796 }
797 }
798 }
799
800 num_nz_ptr[ridx_l] = icol;
801 });
802 }
804 }
805 }
806
807 nnz_actual = std::max(nnz_actual, *(nnz_actual_d.copyToHost()));
808 }
809
811 return (nnz_actual - nnz_max);
812}
813
814template <class T, class Ops>
815void MatrixPC<T,Ops>::Apply (T& a_x, const T& a_b)
816{
817 // Given a right-hand-side b, solve:
818 // A x = b
819 // where A is the linear operator, in this case, the curl-curl
820 // operator:
821 // A x = curl (alpha * curl (x) ) + beta * x
822
823 BL_PROFILE("MatrixPC::Apply()");
824 using namespace amrex;
825
827 IsDefined(),
828 "MatrixPC::Apply() called on undefined object" );
830 a_x.getArrayVecType()==warpx::fields::FieldType::Efield_fp,
831 "MatrixPC::Apply() - a_x must be Efield_fp type");
833 a_b.getArrayVecType()==warpx::fields::FieldType::Efield_fp,
834 "MatrixPC::Apply() - a_b must be Efield_fp type");
835
836 WARPX_ABORT_WITH_MESSAGE("MatrixPC<T,Ops>::Apply() - native matrix solvers not implemented. Use with external library, eg, PETSc.");
837
838}
839
840#endif
#define BL_PROFILE(a)
#define AMREX_ALWAYS_ASSERT(EX)
#define AMREX_FORCE_INLINE
#define AMREX_GPU_DEVICE
#define AMREX_GPU_HOST_DEVICE
amrex::ParmParse pp
#define AMREX_D_DECL(a, b, c)
#define WARPX_ABORT_WITH_MESSAGE(MSG)
Definition TextMsg.H:15
#define WARPX_ALWAYS_ASSERT_WITH_MESSAGE(EX, MSG)
Definition TextMsg.H:13
bool IsDefined() const override
Check if the nonlinear solver has been defined.
Definition MatrixPC.H:171
void Apply(T &, const T &) override
Apply (solve) the preconditioner given a RHS.
Definition MatrixPC.H:815
void Define(const T &, Ops *) override
Define the preconditioner.
Definition MatrixPC.H:229
Ops * m_ops
Definition MatrixPC.H:180
bool m_verbose
Definition MatrixPC.H:178
amrex::Gpu::DeviceVector< int > m_r_indices_g
Definition MatrixPC.H:193
MatrixPC(const MatrixPC &)=delete
amrex::Gpu::DeviceVector< int > m_c_indices_g
Definition MatrixPC.H:195
void printParameters() const override
Print parameters.
Definition MatrixPC.H:212
void Update(const T &a_U) override
Update the preconditioner.
Definition MatrixPC.H:284
amrex::Gpu::DeviceVector< int > m_num_nz
Definition MatrixPC.H:194
void setName(const std::string &a_name) override
Set the name for screen output and parsing inputs.
Definition MatrixPC.H:173
void getPCMatrix(amrex::Gpu::DeviceVector< int > &a_r_indices_g, amrex::Gpu::DeviceVector< int > &a_num_nz, amrex::Gpu::DeviceVector< int > &a_c_indices_g, amrex::Gpu::DeviceVector< RT > &a_a_ij, int &a_n, int &a_ncols_max) override
Get the sparse matrix form of the preconditioner.
Definition MatrixPC.H:128
bool m_pc_diag_only
Definition MatrixPC.H:187
const amrex::Vector< amrex::Array< amrex::MultiFab *, 3 > > * m_bcoefs
Definition MatrixPC.H:198
int m_ndofs_l
Definition MatrixPC.H:185
int m_num_realloc
Definition MatrixPC.H:200
MatrixPC()=default
Default constructor.
amrex::Vector< amrex::Geometry > m_geom
Definition MatrixPC.H:183
int m_num_amr_levels
Definition MatrixPC.H:182
typename T::value_type RT
Definition MatrixPC.H:80
void readParameters()
Read parameters.
Definition MatrixPC.H:221
MatrixPC & operator=(const MatrixPC &)=delete
int m_ndofs_g
Definition MatrixPC.H:186
bool m_include_mass_matrices
Definition MatrixPC.H:189
int Assemble(const T &a_U)
Assemble the matrix.
Definition MatrixPC.H:317
int m_pc_mat_nnz
Definition MatrixPC.H:188
std::string m_name
Definition MatrixPC.H:191
~MatrixPC() override=default
Default destructor.
MatrixPC(MatrixPC &&) noexcept=delete
bool m_is_defined
Definition MatrixPC.H:177
amrex::Gpu::DeviceVector< amrex::Real > m_a_ij
Definition MatrixPC.H:196
RT m_dt
Definition Preconditioner.H:117
Preconditioner()=default
Default constructor.
__host__ __device__ bool contains(const IntVectND< dim > &p) const noexcept
Array4< typename FabArray< FAB >::value_type const > const_array(const MFIter &mfi) const noexcept
T const * data() const noexcept
amrex_real Real
PODVector< T, ArenaAllocator< T > > DeviceVector
Definition MatrixPC.H:24
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE bool insertOrAdd(const int a_cidx, const amrex::Real a_val, int *const a_cidxs, amrex::Real *const a_aij, const int a_nnz, int &a_ncol)
Definition MatrixPC.H:26
constexpr auto c
vacuum speed of light [m/s]
Definition constant.H:153
void WMRecordWarning(const std::string &topic, const std::string &text, const WarnPriority &priority=WarnPriority::medium)
Helper function to abbreviate the call to WarnManager::GetInstance().RecordWarning (recording a warni...
Definition WarnManager.cpp:320
__host__ __device__ AMREX_FORCE_INLINE T Max(T *const m, T const value) noexcept
void synchronize() noexcept
void copyAsync(HostToDevice, InIter begin, InIter end, OutIter result) noexcept
static constexpr DeviceToDevice deviceToDevice
void streamSynchronize() noexcept
__host__ __device__ void ignore_unused(const Ts &...)
std::enable_if_t< std::is_integral_v< T > > ParallelFor(TypeList< CTOs... > ctos, std::array< int, sizeof...(CTOs)> const &runtime_options, T N, F &&f)
BoxND< 3 > Box
IntVectND< 3 > IntVect
@ Efield_fp
Definition Fields.H:97