Compare commits

..
Author SHA1 Message Date
Will Pazner 600b30c661 Add more annotations, modifying LOR solvers miniapp for benchmarking 2025-07-10 20:39:40 -07:00
Will Pazner ad9ba9d1ba Add a couple more perf function annotations 2025-07-10 15:26:57 -07:00
Will Pazner 3d2f81530a Merge remote-tracking branch 'origin/ceed-solver-bps' into ceed-bp-bps-tests 2025-07-09 09:34:29 -07:00
Veselin Dobrev 539f958a9b Fix merge issue 2025-07-09 01:01:59 -07:00
Veselin Dobrev a0b2a77b69 Merge branch 'master' into ceed-solver-bps 2025-07-09 00:50:01 -07:00
Veselin Dobrev e2292a93b8 MFEM_PERF_* annotations: added options to enable GPU-stream- and
MPI-synchronizations at the start and at the end of annotation regions. These
synchronizations can be enabled or disabled (default) in code via the new
macros: MFEM_PERF_SYNC, MFEM_PERF_SYNC_STREAM, and MFEM_PERF_SYNC_MPI; the
environment variables with the same names can be set to 0/1 to control the
synchronization as well.
2025-07-08 19:23:09 -07:00
Will Pazner 22384799dd Add some Caliper annotations 2025-07-08 10:33:22 -07:00
Veselin Dobrev a5943b33d3 Merge branch 'pbench_ceed' into ceed-bp-bps-tests 2025-07-07 18:47:16 -07:00
Veselin Dobrev 9821670831 In pbench_ceed.cpp, fix the BK* tests to use E-to-E vector action 2025-07-07 18:44:40 -07:00
Veselin Dobrev d13a8bfa22 Merge branch 'device-conf-prolongation-optim' into ceed-bp-bps-tests 2025-07-06 10:57:47 -07:00
Veselin Dobrev e38eeccc2e Merge branch 'pbench_ceed' into ceed-bp-bps-tests 2025-07-06 10:29:52 -07:00
Veselin Dobrev fce5095ba0 Merge branch 'master' into ceed-solver-bps
Resolved conflicts:
   .gitignore
   makefile
2025-07-05 14:00:06 -07:00
Veselin Dobrev 408c8e9468 Redefine MFEM_MPI_CXX_BOOL as MPI_UNSIGNED_CHAR because in some
implementations MFEM_MPI_CXX_BOOL does not work correctly.
2025-07-05 13:56:13 -07:00
Veselin Dobrev 1b56b05b8d Merge branch 'master' into pbench_ceed 2025-07-05 13:54:39 -07:00
Veselin Dobrev cc279f9ea5 Redefine MFEM_MPI_CXX_BOOL as MPI_UNSIGNED_CHAR because in some
implementations MFEM_MPI_CXX_BOOL does not work correctly.
2025-07-05 13:52:29 -07:00
Veselin Dobrev b8eb1ecec8 Small tweaks 2025-05-19 22:23:35 -07:00
Veselin Dobrev fe679e0e05 Add a parallel version of the CEED benchmarks, pbench_ceed.cpp, in
tests/benchmarks.
2025-05-13 16:31:04 -07:00
Veselin Dobrev 602e889600 Add miniapps/benchmarks/ceed-solver-bps/solver-bp to .gitignore 2025-04-20 14:48:42 -07:00
Veselin Dobrev 4a82e0112f Fix the single precision build in miniapps/benchmarks/ceed-solver-bps 2025-04-19 16:33:46 -07:00
Veselin Dobrev 198c35a314 Merge branch 'master' into ceed-solver-bps 2025-04-19 16:20:57 -07:00
Veselin Dobrev bd50b287aa Fix style 2025-04-19 16:04:55 -07:00
Veselin Dobrev 4244d97e1f Fix GPU and out-of-source builds for miniapps/benchmarks/ceed-solver-bps 2025-04-19 15:59:00 -07:00
Veselin Dobrev 75a4018d9c Initial version of the ceed-solver-bps benchmark in miniapps/benchmarks 2025-04-17 16:10:20 -07:00
Veselin Dobrev 72323aad24 Merge branch 'master' into ceed-solver-bps
Resolved conflicts:
  fem/transfer.hpp
2025-04-17 15:45:51 -07:00
Veselin Dobrev ce9f969f16 Some tweaks in class MultigridBase to avoid implicit conversions of
call parameters.

Some small optimizations in class OperatorChebyshevSmoother.
2025-04-17 15:37:54 -07:00
Veselin Dobrev 6e7718db30 Optimizations for class DeviceConformingProlongationOperator 2025-04-03 10:07:28 -07:00
Veselin Dobrev 1fe1e50598 Merge branch 'mg-smem-prolongation' into ceed-solver-bps 2025-02-20 23:26:31 -08:00
Veselin Dobrev f0e542787b Fix memory leaks in miniapps/solvers/block-solvers.cpp 2025-02-20 23:24:35 -08:00
Veselin Dobrev 9627443feb Fix uninitialized memory usage due to incorrect setting of the
iterative mode of the MINRESSolver in BDPMinresSolver in
miniapps/solvers/darcy_solver.?pp. This became an issue due to the way
MultigridBase::SmoothingStep now calls the smoothers (including the
coarse solver) -- it sets the iterative mode of the smoother to false
and calls its Mult() method without initializing the second argument.
2025-02-20 21:48:23 -08:00
Veselin Dobrev ac2252ff20 In class MultigridBase, added support for coarse solver that is
used in addition to the coarse smoother, if given.

In class OperatorChebyshevSmoother, added support for changing the
order after construction with a new method: SetOrder().

Added a few minor tweaks.
2025-02-20 17:17:03 -08:00
Veselin Dobrev 3dc1217b30 Merge branch 'solver-monitor-controller' into ceed-solver-bps 2025-02-20 13:49:14 -08:00
Veselin Dobrev 32c5af1151 Merge branch 'mg-smem-prolongation' into ceed-solver-bps 2025-02-20 13:47:10 -08:00
Veselin Dobrev f573b345ae double -> real_t 2025-02-19 18:16:53 -08:00
Veselin Dobrev 80aaf8bb0d Merge branch 'master' into mg-smem-prolongation
Resolved conflicts:
   examples/ex26.cpp
   examples/ex26p.cpp
   fem/transfer.cpp
   linalg/solvers.cpp
2025-02-19 18:11:09 -08:00
Veselin Dobrev a8ea2cf89a In MultigridBase::Cycle, fix for the case of no pre-smoothing 2024-02-09 22:28:42 -08:00
Veselin Dobrev 5768645171 In class MultigridBase:
* add support for iterative mode
* avoid unnecessary initializations with zero and residual computations
  when using W-cycle
* ensure smoothers (including the level 0 coarse solver) are used only
  with iterative mode turned off
* when using W-cycle, do not call the level 0  coarse solver twice to
  avoid calling "exact" solvers twice.
2024-02-09 21:13:02 -08:00
Veselin Dobrev c2a721cabe Bugfix: add missing allocation of OperatorChebyshevSmoother::z 2024-02-02 06:45:00 -08:00
Veselin Dobrev 343ce0c26d Fix the arguments in calls to SmemProlongationTranspose3D 2024-02-02 06:17:53 -08:00
Veselin Dobrev 5c4d341621 Fix a typo from the previous commit: 'R_[j]' -> '*R_[j]'.
Rename Operator::z to Operator::z_am to avoid confusion with the
Vectors z defined in a number of derived classes, and make it
private for reasons explained in the added doxygen comment.

Note: this commit does not seem to pass some of the unit tests.
2024-02-02 05:14:57 -08:00
Veselin Dobrev 35c9c43012 In OperatorChebyshevSmoother::Mult, for k = 0, use a single kernel
instead of 3 separate kernels.
2024-02-02 04:38:37 -08:00
Veselin Dobrev 371b091fd1 Use 'subtract' in a few places instead of the more general 'add' 2024-02-02 03:58:35 -08:00
Veselin Dobrev 1ed1ba8a71 Add optimization for OperatorJacobiSmoother::Mult from PR #1852 2024-02-02 03:41:58 -08:00
Will Pazner 2bd73092d4 Eliminate two more copies in multigrid 2024-02-01 15:26:38 -08:00
Will Pazner aeb4d508bb Remove temporaries from MultigridBase
Add temporary vector z in Operator
2024-02-01 15:26:24 -08:00
Will Pazner a58d941918 Small optimization in MultigridBase::Cycle (requires coarse solver not in iterative mode) 2024-02-01 15:16:45 -08:00
Will Pazner 60a7a9e673 Reduce copies in MultigridBase::Cycle and MultigridBase::SmoothingStep 2024-02-01 15:16:45 -08:00
Will Pazner a1d643f287 Small improvements to OperatorChebyshevSmoother
- Fuse two kernels in OperatorChebyshevSmoother::Mult
- Remove unnecessary copy
- Improve unit test
2024-02-01 14:47:06 -08:00
Will Pazner 04f303fd09 Add multigrid prolongation specializations 2024-02-01 12:04:17 -08:00
Will Pazner 1c4390c0f1 Add optimized p-multigrid prolongation kernels
Also rename Restriction to ProlongationTranspose
2024-02-01 10:49:07 -08:00
Will Pazner 35121d7297 Remove TensorProductPRefinementTransferOperator::~TensorProductPRefinementTransferOperator 2024-02-01 10:48:09 -08:00
126 changed files with 4093 additions and 4063 deletions
+5 -8
View File
@@ -168,13 +168,10 @@ jobs:
env
shell: bash
# For info on Xcode see:
# - https://github.com/actions/runner-images/issues/12541
# - https://github.com/actions/runner-images/blob/releases/macos-15-arm64/20250811/images/macos/macos-15-arm64-Readme.md#xcode
- name: Xcode version setup (MacOS)
if: matrix.os == 'macos-latest'
run: |
XCODE_PATH="/Applications/Xcode_16.4.app"
XCODE_PATH="/Applications/Xcode_15.3.app"
echo "> sudo xcode-select -s ${XCODE_PATH}"
sudo xcode-select -s ${XCODE_PATH}
echo "> g++ -v"
@@ -292,10 +289,10 @@ jobs:
run: |
export HOMEBREW_NO_INSTALL_CLEANUP=1
brew update
brew install llvm@20 enzyme
echo "LLVM_PREFIX=$(brew --prefix llvm@20)" >> $GITHUB_ENV
echo "OMPI_CC=$(brew --prefix llvm@20)/bin/clang" >> $GITHUB_ENV
echo "OMPI_CXX=$(brew --prefix llvm@20)/bin/clang++" >> $GITHUB_ENV
brew install llvm@19 enzyme
echo "LLVM_PREFIX=$(brew --prefix llvm@19)" >> $GITHUB_ENV
echo "OMPI_CC=$(brew --prefix llvm@19)/bin/clang" >> $GITHUB_ENV
echo "OMPI_CXX=$(brew --prefix llvm@19)/bin/clang++" >> $GITHUB_ENV
# MFEM build and test
- name: build
+7 -6
View File
@@ -211,7 +211,7 @@ miniapps/electromagnetics/joule
miniapps/electromagnetics/Volta-AMR*
miniapps/electromagnetics/Tesla-AMR*
miniapps/electromagnetics/Maxwell-Parallel*
miniapps/electromagnetics/Joule_[0-9]*
miniapps/electromagnetics/Joule_*
miniapps/gslib/field-diff
miniapps/gslib/field-interp
@@ -267,9 +267,9 @@ miniapps/meshing/bounding-box*
miniapps/meshing/jacobian-determinant*
miniapps/mtop/parheat
miniapps/mtop/ParHeat/*
miniapps/mtop/ParHeat*
miniapps/mtop/seqheat
miniapps/mtop/SeqHeat/*
miniapps/mtop/SeqHeat*
miniapps/autodiff/paradiff
miniapps/autodiff/seqadiff
@@ -277,7 +277,7 @@ miniapps/autodiff/seqtest
miniapps/autodiff/par_example
miniapps/autodiff/seq_example
miniapps/autodiff/seq_test
miniapps/autodiff/Example/*
miniapps/autodiff/Exampl*
miniapps/navier/navier_mms
miniapps/navier/navier_kovasznay
@@ -300,7 +300,6 @@ miniapps/nurbs/nurbs_solenoidal
miniapps/nurbs/nurbs_printfunc
miniapps/nurbs/nurbs_patch_ex1
miniapps/nurbs/nurbs_curveint
miniapps/nurbs/nurbs_surface
miniapps/nurbs/refined.mesh
miniapps/nurbs/mesh.*
miniapps/nurbs/sol_?.gf
@@ -319,7 +318,6 @@ miniapps/nurbs/nurbs_naca_cmesh
miniapps/nurbs/naca-cmesh.mesh
miniapps/nurbs/glvis_naca-cmesh.mesh
miniapps/nurbs/Naca_cmesh
miniapps/nurbs/*-Surface.mesh
miniapps/performance/ex1
miniapps/performance/ex1p
@@ -411,6 +409,8 @@ miniapps/tribol/contact-patch-test
miniapps/diag-smoothers/abs-l1-jacobi
miniapps/diag-smoothers/mg-abs-l1-jacobi
miniapps/benchmarks/ceed-solver-bps/solver-bp
# Unit test binary and outputs
tests/unit/output_meshes
tests/unit/unit_tests
@@ -431,6 +431,7 @@ tests/benchmarks/bench_ceed
tests/benchmarks/bench_tmop
tests/benchmarks/bench_vector
tests/benchmarks/bench_virtuals
tests/benchmarks/pbench_ceed
# Test script output
tests/scripts/*.err
+5 -5
View File
@@ -22,7 +22,7 @@ include:
# the "needs" keyword and express the DAG of jobs for more efficiency.
# - We use setup and setup_baseline phases to download content outside of mfem
# directory.
# - Allocate/Release is where Dane resource are allocated/released once for all.
# - Allocate/Release is where ruby resource are allocated/released once for all.
# - Build and Test is where we build and MFEM for multiple toolchains.
# - Baseline_checks gathers baseline-type test suites execution
# - Baseline_publish, only available on master, allows to update baseline
@@ -53,7 +53,7 @@ variables:
AUTOTEST_COMMIT: "YES"
# Trigger subpipelines:
dane-build-and-test:
ruby-build-and-test:
stage: sub-pipelines
variables:
# Explicitly pass down values that we want to be able to set when triggering
@@ -61,10 +61,10 @@ dane-build-and-test:
AUTOTEST: "${AUTOTEST}"
AUTOTEST_COMMIT: "${AUTOTEST_COMMIT}"
trigger:
include: .gitlab/dane-build-and-test.yml
include: .gitlab/ruby-build-and-test.yml
strategy: depend
dane-baseline:
ruby-baseline:
stage: sub-pipelines
variables:
# Explicitly pass down values that we want to be able to set when triggering
@@ -73,7 +73,7 @@ dane-baseline:
AUTOTEST: "${AUTOTEST}"
AUTOTEST_COMMIT: "${AUTOTEST_COMMIT}"
trigger:
include: .gitlab/dane-baseline.yml
include: .gitlab/ruby-baseline.yml
strategy: depend
lassen-build-and-test:
+3 -3
View File
@@ -24,7 +24,7 @@ and `test type`.
Machines typically include:
* Dane: Intel Sapphire Rapids
* Ruby: 2nd Gen Intel Xeon (Cascade Lake)
* Lassen: Power9 + Nvidia GPU
* Corona: AMD GPU
@@ -76,13 +76,13 @@ with a spack spec of MFEM, within the limits permitted by the MFEM spack
package.
In any build-and-test sub-pipeline a job basically consists in defining the
spack spec to use. Adding a job on Dane for example resumes to:
spack spec to use. Adding a job on ruby for example resumes to:
```yaml
<job_name>:
variables:
SPEC: "<spack_spec>"
extends: .build_and_test_on_dane
extends: .build_and_test_on_ruby
```
The remaining and non trivial work is to make sure this spec is working. To
+1 -1
View File
@@ -24,7 +24,7 @@ variables:
# TODO: add a clean-up mechanism
BUILD_ROOT: ${USER_CI_TOP_DIR}/${CI_PROJECT_NAME}-${MACHINE_NAME}-pipeline-${CI_PIPELINE_ID}
# On LLNL's Dane, there is only one allocation shared among jobs in order to
# On LLNL's ruby, there is only one allocation shared among jobs in order to
# save time and resource. This allocation has to be uniquely named so that we
# are sure to retrieve it.
ALLOC_NAME: ${CI_PROJECT_NAME}_ci_${CI_PIPELINE_ID}
@@ -9,17 +9,17 @@
# terms of the BSD-3 license. We welcome feedback and contributions, see file
# CONTRIBUTING.md for details.
# GitLab pipelines configurations for the Dane machine at LLNL
# GitLab pipelines configurations for the Ruby machine at LLNL
variables:
MACHINE_NAME: dane
MACHINE_NAME: ruby
.on_dane:
.on_ruby:
tags:
- shell
- dane
- ruby
rules:
# Don't run dane jobs if...
- if: '$CI_COMMIT_BRANCH =~ /_qnone/ || $ON_DANE == "OFF"'
# Don't run ruby jobs if...
- if: '$CI_COMMIT_BRANCH =~ /_qnone/ || $ON_RUBY == "OFF"'
when: never
# Don't run autotest update if...
- if: '$CI_JOB_NAME =~ /report/ && $AUTOTEST != "YES"'
@@ -40,17 +40,16 @@ variables:
- when: on_success
# Spack helped builds
# Generic dane build job, extending build script
.build_and_test_on_dane:
extends: [.on_dane]
# Generic ruby build job, extending build script
.build_and_test_on_ruby:
extends: [.on_ruby]
stage: build_and_test
script:
# THREADS is used by 'tests/gitlab/build_and_test', run below
# Dane has 224 threads/node and we run 7 separate jobs: 224=7*32
- export THREADS=28
- export THREADS=16
- echo ${ALLOC_NAME}
- export JOBID=$(squeue -h --name=${ALLOC_NAME} --format=%A)
- echo ${JOBID}
- echo ${MFEM_DATA_DIR}
- echo ${SPEC}
- srun $( [[ -n "${JOBID}" ]] && echo "--jobid=${JOBID}" ) --reservation=ci -t 60 -N 1 tests/gitlab/build_and_test --spec "${SPEC}" --data-dir "${MFEM_DATA_DIR}" --data
- srun $( [[ -n "${JOBID}" ]] && echo "--jobid=${JOBID}" ) --reservation=ci -t 45 -N 1 tests/gitlab/build_and_test --spec "${SPEC}" --data-dir "${MFEM_DATA_DIR}" --data
+1 -1
View File
@@ -18,7 +18,7 @@
setup_baseline:
tags:
- shell
- dane
- ruby
stage: setup
variables:
GIT_STRATEGY: none
+1 -1
View File
@@ -16,7 +16,7 @@
setup:
tags:
- shell
- dane
- ruby
stage: setup
variables:
GIT_STRATEGY: none
@@ -19,8 +19,8 @@ stages:
- cleanup
- baseline_publish
baselinecheck_mfem_intel_dane:
extends: [.on_dane]
baselinecheck_mfem_intel_ruby:
extends: [.on_ruby]
stage: baseline_check
variables:
# TPLS_DIR is used in .gitlab/scripts/baseline to provide the tpls location
@@ -31,8 +31,8 @@ baselinecheck_mfem_intel_dane:
script:
- echo ${BUILD_ROOT}
- echo ${TPLS_DIR}
# Used by the tests in MFEM/tests, dane has 224 threads/node:
- export MFEM_TEST_NP=192
# Used by the tests in MFEM/tests:
- export MFEM_TEST_NP=48
# The next script uses the following environment variables:
# * BASELINE_TEST, SYS_TYPE, CI_PROJECT_DIR, ARTIFACTS_DIR,
# * BUILD_ROOT, TPLS_DIR, MACHINE_NAME
@@ -44,7 +44,7 @@ baselinecheck_mfem_intel_dane:
allow_failure: true
cleanup:
extends: .on_dane
extends: .on_ruby
stage: cleanup
variables:
GIT_STRATEGY: none
@@ -53,7 +53,7 @@ cleanup:
- rm -rf "${BUILD_ROOT}" || true
report_baseline:
extends: [.on_dane]
extends: [.on_ruby]
stage: baseline_report
script:
- echo ${MACHINE_NAME}
@@ -113,8 +113,8 @@ report_baseline:
exit $err
) 9> autotest.lock
baselinepublish_mfem_dane:
extends: [.on_dane]
baselinepublish_mfem_ruby:
extends: [.on_ruby]
stage: baseline_publish
rules:
# - if: '$CI_COMMIT_BRANCH == "master" || $REBASELINE == "YES"'
@@ -129,5 +129,5 @@ baselinepublish_mfem_dane:
include:
- local: .gitlab/configs/common.yml
- local: .gitlab/configs/dane-config.yml
- local: .gitlab/configs/ruby-config.yml
- local: .gitlab/configs/setup-baseline.yml
@@ -19,54 +19,54 @@ stages:
allocate_resource:
variables:
GIT_STRATEGY: none
extends: .on_dane
extends: .on_ruby
stage: allocate_resource
script:
- echo ${ALLOC_NAME}
- salloc --exclusive --nodes=1 --reservation=ci --time=60 --no-shell --job-name=${ALLOC_NAME}
timeout: 6h
# GitLab jobs for the Dane machine at LLNL
# GitLab jobs for the Ruby machine at LLNL
debug_ser_gcc_10:
variables:
SPEC: "%gcc@10.3.1 +debug~mpi"
extends: .build_and_test_on_dane
extends: .build_and_test_on_ruby
debug_par_gcc_10:
variables:
SPEC: "%gcc@10.3.1 +debug+mpi"
extends: .build_and_test_on_dane
extends: .build_and_test_on_ruby
opt_ser_gcc_10:
variables:
SPEC: "%gcc@10.3.1 ~mpi"
extends: .build_and_test_on_dane
extends: .build_and_test_on_ruby
opt_par_gcc_10:
variables:
SPEC: "%gcc@10.3.1"
extends: .build_and_test_on_dane
extends: .build_and_test_on_ruby
opt_par_gcc_10_sundials:
variables:
SPEC: "%gcc@10.3.1 +sundials"
extends: .build_and_test_on_dane
extends: .build_and_test_on_ruby
opt_par_gcc_10_petsc:
variables:
SPEC: "%gcc@10.3.1 +petsc ^petsc+mumps~superlu-dist"
extends: .build_and_test_on_dane
extends: .build_and_test_on_ruby
opt_par_gcc_10_pumi:
variables:
SPEC: "%gcc@10.3.1 +pumi"
extends: .build_and_test_on_dane
extends: .build_and_test_on_ruby
# Release
release_resource:
variables:
GIT_STRATEGY: none
extends: .on_dane
extends: .on_ruby
stage: release_resource_and_report
script:
- echo ${ALLOC_NAME}
@@ -78,17 +78,17 @@ release_resource:
report_job_success:
stage: release_resource_and_report
extends:
- .on_dane
- .on_ruby
- .report_job_success
report_job_failure:
stage: release_resource_and_report
extends:
- .on_dane
- .on_ruby
- .report_job_failure
include:
- local: .gitlab/configs/common.yml
- local: .gitlab/configs/dane-config.yml
- local: .gitlab/configs/ruby-config.yml
- local: .gitlab/configs/setup-build-and-test.yml
- local: .gitlab/configs/report-build-and-test.yml
+2 -2
View File
@@ -14,7 +14,7 @@
# locals
glob_err=${BASELINE_TEST}.err
base=${BASELINE_TEST}-${SYS_TYPE}
if [[ "${MACHINE_NAME}" == "dane" ]]; then
if [[ "${MACHINE_NAME}" == "ruby" ]]; then
base="${BASELINE_TEST}-${MACHINE_NAME}"
fi
base_diff=${base}.diff
@@ -31,7 +31,7 @@ cd tests
mkdir _${BASELINE_TEST} && cd _${BASELINE_TEST}
# run
if [[ "${MACHINE_NAME}" == "dane" ]]; then
if [[ "${MACHINE_NAME}" == "ruby" ]]; then
salloc --nodes=1 --exclusive --reservation=ci ../runtest ../../mfem "${BASELINE_TEST} ${TPLS_DIR}"
elif [[ ${MACHINE_NAME} == "corona" ]]; then
salloc --nodes=1 -t 60 -p pbatch ../runtest ../../mfem "${BASELINE_TEST} ${TPLS_DIR}"
+2 -2
View File
@@ -11,7 +11,7 @@
# terms of the BSD-3 license. We welcome feedback and contributions, see file
# CONTRIBUTING.md for details.
# There will be collision between corona and dane baselines.
# There will be collision between corona and ruby baselines.
# Once the corresponding files have been generated, we can switch to machine
# specific ref.
ARTIFACT_PATH=${CI_PROJECT_DIR}/${ARTIFACTS_DIR}/baseline-${SYS_TYPE}
@@ -21,7 +21,7 @@ PATCH_FILE=${ARTIFACT_PATH}.patch
FULL_FILE=${ARTIFACT_PATH}.out
DIFF_FILE=${ARTIFACT_PATH}.diff
# There will be collision between corona and dane baselines.
# There will be collision between corona and ruby baselines.
# Once the corresponding files have been generated, we can switch to machine
# specific ref.
SAVED_NAME=baseline-${SYS_TYPE}.saved
+7 -10
View File
@@ -29,14 +29,9 @@ Discretization improvements
Meshing improvements
--------------------
- Added support for higher order meshes in Mesh::MakeSimplicial and
ParMesh::MakeSimplicial.
- Added a new miniapp for interpolating a surface grid of points in 3D using a
smooth NURBS surface, that can then be sampled at arbitrary resolution while
staying close to the original geometry. See miniapps/nurbs/nurbs_surface.
GPU computing
-------------
- The function Vector::SetSubVector(const Array<int> &, const real_t) now
@@ -46,11 +41,6 @@ GPU computing
where host execution is always needed (e.g. when the DOFs array is small).
- Introduced MFEM_FOREACH_THREAD_DIRECT, which directly maps loop tasks to GPU
threads, assigning one task per thread.
- Implemented a GPU-accelerated matrix-free AMR derefinement `GridFunction`
update operator. This supports mixed geometry meshes and variable order
spaces, and is the default derefinement operator constructed by
`FiniteElementSpace::Update` and `ParFiniteElementSpace::Update`.
The operator requires `FiniteElementSpace::Nonconforming() == true`.
New and updated examples and miniapps
-------------------------------------
@@ -80,6 +70,13 @@ Miscellaneous
variable is an alternative to calling 'Device::SetGPUAwareMPI(true)'.
- Added parallel Address Sanitizer, serial and parallel Undefined Behavior
Sanitizer and serial Memory Sanitizer GitHub actions tests on Ubuntu.
- MFEM_PERF_* annotations: added options to enable GPU-stream- and
MPI-synchronizations at the start and at the end of annotation regions. These
synchronizations can be enabled or disabled (default) in code via the new
macros: MFEM_PERF_SYNC, MFEM_PERF_SYNC_STREAM, and MFEM_PERF_SYNC_MPI; the
environment variables with the same names can be set to 0/1 to control the
synchronization as well.
Version 4.8, released on Apr 9, 2025
====================================
+5 -14
View File
@@ -598,20 +598,14 @@ set(MFEM_TPLS OPENMP HYPRE LAPACK BLAS SuperLUDist STRUMPACK METIS SuiteSparse
BENCHMARK PARELAG TRIBOL MPI_CXX HIP HIPBLAS HIPSPARSE MOONOLITH BLITZ
ALGOIM ENZYME)
# Add all created targets and *_FOUND libraries in the variables TPL_TARGETS and
# TPL_LIBRARIES, respectively.
set(TPL_TARGETS)
# Add all *_FOUND libraries in the variable TPL_LIBRARIES.
set(TPL_LIBRARIES "")
set(TPL_INCLUDE_DIRS "")
foreach(TPL IN LISTS MFEM_TPLS)
if (${TPL}_FOUND OR TARGET ${TPL})
if (${TPL}_FOUND)
message(STATUS "MFEM: using package ${TPL}")
if (TARGET ${TPL})
list(APPEND TPL_TARGETS ${TPL})
else()
list(APPEND TPL_LIBRARIES ${${TPL}_LIBRARIES})
list(APPEND TPL_INCLUDE_DIRS ${${TPL}_INCLUDE_DIRS})
endif()
list(APPEND TPL_LIBRARIES ${${TPL}_LIBRARIES})
list(APPEND TPL_INCLUDE_DIRS ${${TPL}_INCLUDE_DIRS})
endif()
endforeach(TPL)
list(REVERSE TPL_LIBRARIES)
@@ -686,10 +680,7 @@ set(MFEM_INSTALL_DIR ${CMAKE_INSTALL_PREFIX})
# Declaring the library
mfem_add_library(mfem ${SOURCES} ${HEADERS} ${MASTER_HEADERS})
# message(STATUS "TPL_LIBRARIES = ${TPL_LIBRARIES}")
target_link_libraries(mfem PUBLIC ${TPL_LIBRARIES} ${TPL_TARGETS})
if (TPL_TARGETS)
add_dependencies(mfem ${TPL_TARGETS})
endif()
target_link_libraries(mfem PUBLIC ${TPL_LIBRARIES})
if (MINGW)
target_link_libraries(mfem PRIVATE ws2_32)
endif()
-16
View File
@@ -121,11 +121,6 @@ Parallel build:
make -j 4
(For METIS 5, see https://mfem.org/building/#parallel-build-using-metis-5)
Parallel build with fetching of hypre and METIS:
mkdir <mfem-buil-dir> ; cd <mfem-build-dir>
cmake <mfem-source-dir> -DMFEM_USE_MPI=YES -DFETCH_TPLS=YES
make -j 4
CUDA build:
(this build requires CMake 3.17 or newer)
mkdir <mfem-build-dir> ; cd <mfem-build-dir>
@@ -668,7 +663,6 @@ The specific libraries and their options are:
- OpenMP (optional), usually part of compiler, used when either MFEM_USE_OPENMP
or MFEM_USE_LEGACY_OPENMP is set to YES.
Options: OPENMP_OPT, OPENMP_LIB.
Versions: OpenMP >= 3.1 when MFEM_USE_OPENMP=YES.
- High-resolution POSIX clocks: when using MFEM_TIMER_TYPE = 2, it may be
necessary to link with a system library (e.g. librt.so).
@@ -848,7 +842,6 @@ The specific libraries and their options are:
- HIP (optional), used when MFEM_USE_HIP = YES.
URL: https://rocmdocs.amd.com
Options: HIP_CXX, HIP_ARCH, HIP_OPT, HIP_LIB.
Versions: ROCm >= 5.6.1.
- OCCA (optional), used when MFEM_USE_OCCA = YES.
URL: https://libocca.org
@@ -1081,9 +1074,6 @@ The following options are CMake specific:
MFEM_ENABLE_TESTING - Enable the ctest framework for testing.
MFEM_ENABLE_EXAMPLES - Build all of the examples by default.
MFEM_ENABLE_MINIAPPS - Build all of the miniapps by default.
FETCH_TPLS - Enable fetching of all supported third-party libraries.
HYPRE_FETCH - Enable fetching of hypre.
METIS_FETCH - Enable fetching of metis.
External libraries (CMake):
---------------------------
@@ -1145,12 +1135,6 @@ The following built-in CMake packages are also used:
set the <LIBNAME>_LIBRARIES option directly; the configuration option
<LIBNAME>_DIR is not supported.
The MFEM CMake build system also provides fetching (automated building) for the
packages/libraries listed below. Note that when fetching is enabled, any related
auto-detection functionality is disabled.
- HYPRE
- METIS
Building without GNU make or CMake
==================================
+2 -54
View File
@@ -9,18 +9,15 @@
# terms of the BSD-3 license. We welcome feedback and contributions, see file
# CONTRIBUTING.md for details.
# Defines the following variables if fetching of TPLs is disabled (default):
# Defines the following variables:
# - HYPRE_FOUND
# - HYPRE_LIBRARIES
# - HYPRE_INCLUDE_DIRS
# - HYPRE_VERSION
# - HYPRE_USING_CUDA (internal)
# - HYPRE_USING_HIP (internal)
# otherwise, the following are defined:
# - HYPRE (imported library target)
# - HYPRE_VERSION (cache variable)
if (HYPRE_FOUND OR TARGET HYPRE)
if (HYPRE_FOUND)
if (HYPRE_USING_CUDA)
find_package(CUDAToolkit REQUIRED)
endif()
@@ -36,55 +33,6 @@ if (HYPRE_FOUND OR TARGET HYPRE)
endif()
endif()
if (HYPRE_FETCH OR FETCH_TPLS)
set(HYPRE_FETCH_VERSION 2.33.0)
add_library(HYPRE STATIC IMPORTED)
# set options and associated dependencies
set(CMAKE_OPTIONS)
list(APPEND CMAKE_OPTIONS -DCMAKE_BUILD_TYPE:STRING=${CMAKE_BUILD_TYPE})
if (MFEM_USE_CUDA)
list(APPEND CMAKE_OPTIONS -DHYPRE_WITH_CUDA:BOOL=ON)
find_package(CUDAToolkit REQUIRED)
target_link_libraries(HYPRE INTERFACE CUDA::cusparse CUDA::curand CUDA::cublas)
elseif (MFEM_USE_HIP)
list(APPEND CMAKE_OPTIONS -DHYPRE_WITH_HIP:BOOL=ON)
find_package(rocsparse REQUIRED)
find_package(rocrand REQUIRED)
target_link_libraries(HYPRE INTERFACE rocsparse rocrand)
endif()
if (MFEM_USE_SINGLE)
list(APPEND CMAKE_OPTIONS -DHYPRE_ENABLE_SINGLE:BOOL=ON)
endif()
# define external project and create future include directory so it is present
# to pass CMake checks at end of MFEM configuration step
message(STATUS "Will fetch HYPRE ${HYPRE_FETCH_VERSION} to be built with ${CMAKE_OPTIONS}")
set(PREFIX ${CMAKE_BINARY_DIR}/fetch/hypre)
include(ExternalProject)
ExternalProject_Add(hypre
GIT_REPOSITORY https://github.com/hypre-space/hypre.git
GIT_TAG v${HYPRE_FETCH_VERSION}
GIT_SHALLOW TRUE
UPDATE_DISCONNECTED TRUE
SOURCE_SUBDIR src
PREFIX ${PREFIX}
CMAKE_CACHE_ARGS -DCMAKE_INSTALL_PREFIX:PATH=${PREFIX} -DCMAKE_INSTALL_LIBDIR:PATH=lib ${CMAKE_OPTIONS})
file(MAKE_DIRECTORY ${PREFIX}/include)
# set imported library target properties
add_dependencies(HYPRE hypre)
set_target_properties(HYPRE PROPERTIES
IMPORTED_LOCATION ${PREFIX}/lib/libHYPRE.a
INTERFACE_INCLUDE_DIRECTORIES ${PREFIX}/include)
# convert HYPRE version to integer
string(REGEX MATCHALL "[0-9]+" HYPRE_SPLIT_VERSION ${HYPRE_FETCH_VERSION})
list(GET HYPRE_SPLIT_VERSION 0 HYPRE_MAJOR_VERSION)
list(GET HYPRE_SPLIT_VERSION 1 HYPRE_MINOR_VERSION)
list(GET HYPRE_SPLIT_VERSION 2 HYPRE_PATCH_VERSION)
math(EXPR HYPRE_VERSION "10000*${HYPRE_MAJOR_VERSION} + 100*${HYPRE_MINOR_VERSION} + ${HYPRE_PATCH_VERSION}")
# set cache variables that would otherwise be set after mfem_find_package call
set(HYPRE_VERSION ${HYPRE_VERSION} CACHE STRING "HYPRE version." FORCE)
return()
endif()
include(MfemCmakeUtilities)
mfem_find_package(HYPRE HYPRE HYPRE_DIR "include" "HYPRE.h" "lib" "HYPRE"
"Paths to headers required by HYPRE." "Libraries required by HYPRE."
+1 -29
View File
@@ -9,38 +9,10 @@
# terms of the BSD-3 license. We welcome feedback and contributions, see file
# CONTRIBUTING.md for details.
# Defines the following variables if fetching of TPLs is disabled (default):
# Defines the following variables:
# - METIS_FOUND
# - METIS_LIBRARIES
# - METIS_INCLUDE_DIRS
# - METIS_VERSION_5
# otherwise, the following are defined:
# - METIS (imported library target)
# - METIS_VERSION_5 (cache variable)
if (METIS_FETCH OR FETCH_TPLS)
set(METIS_FETCH_VERSION 4.0.3)
add_library(METIS STATIC IMPORTED)
# define external project
message(STATUS "Will fetch METIS ${METIS_FETCH_VERSION} to be built with default options")
set(PREFIX ${CMAKE_BINARY_DIR}/fetch/metis)
include(ExternalProject)
ExternalProject_Add(metis
GIT_REPOSITORY https://github.com/mfem/tpls
GIT_TAG b60352fbe9675d374b00828055e55be4584c7995 # tag from 1/16/25
GIT_SHALLOW TRUE
UPDATE_DISCONNECTED TRUE
PREFIX ${PREFIX}
CONFIGURE_COMMAND tar -xzf ../metis/metis-${METIS_FETCH_VERSION}-mac.tgz --strip=1
INSTALL_COMMAND mkdir -p ${PREFIX}/lib && cp libmetis.a ${PREFIX}/lib/)
# set imported library target properties
add_dependencies(METIS metis)
set_target_properties(METIS PROPERTIES
IMPORTED_LOCATION ${PREFIX}/lib/libmetis.a)
# set cache variables that would otherwise be set after mfem_find_package call
set(METIS_VERSION_5 FALSE CACHE BOOL "Is METIS version 5?")
return()
endif()
include(MfemCmakeUtilities)
mfem_find_package(METIS METIS METIS_DIR "include;Lib" "metis.h"
+1 -4
View File
@@ -23,14 +23,11 @@
#include "_config.hpp"
#endif
#include <cstdint>
#include <climits>
namespace mfem
{
#if (defined(MFEM_USE_CUDA) && defined(__CUDACC__)) || \
(defined(MFEM_USE_HIP) && defined(__HIP__))
(defined(MFEM_USE_HIP) && defined(__HIPCC__))
#define MFEM_HOST_DEVICE __host__ __device__
#else
#define MFEM_HOST_DEVICE
-6
View File
@@ -89,12 +89,6 @@ option(MFEM_ENABLE_EXAMPLES "Build all of the examples" OFF)
option(MFEM_ENABLE_MINIAPPS "Build all of the miniapps" OFF)
option(MFEM_ENABLE_BENCHMARKS "Build all of the benchmarks" OFF)
# Allow a user to specify fetching of certain third-party libraries instead of
# searching for existing installations.
option(FETCH_TPLS "Enable fetching of all supported third-party libraries" OFF)
option(HYPRE_FETCH "Enable fetching of hypre" OFF)
option(METIS_FETCH "Enable fetching of METIS" OFF)
# Setting CXX/MPICXX on the command line or in user.cmake will overwrite the
# autodetected C++ compiler.
# set(CXX g++)
+3 -1
View File
@@ -78,6 +78,7 @@ private:
opr.SetOperatorOwner(false);
CGSolver* pcg = new CGSolver();
// pcg->iterative_mode = false; // the multigrid algorithm does this
pcg->SetPrintLevel(-1);
pcg->SetMaxIter(200);
pcg->SetRelTol(sqrt(1e-4));
@@ -100,7 +101,8 @@ private:
Vector diag(fespace.GetTrueVSize());
bfs[level]->AssembleDiagonal(diag);
Solver* smoother = new OperatorChebyshevSmoother(*opr, diag, ess_tdof_list, 2);
Solver *smoother = new OperatorChebyshevSmoother(
*opr, diag, ess_tdof_list, 2);
AddLevel(opr.Ptr(), smoother, true, true);
}
};
+1
View File
@@ -88,6 +88,7 @@ private:
amg->SetPrintLevel(-1);
CGSolver* pcg = new CGSolver(MPI_COMM_WORLD);
// pcg->iterative_mode = false; // the multigrid algorithm does this
pcg->SetPrintLevel(-1);
pcg->SetMaxIter(10);
pcg->SetRelTol(sqrt(1e-4));
-2
View File
@@ -82,8 +82,6 @@ set(SRCS
fe/fe_ser.cpp
fe_coll.cpp
fespace.cpp
derefmat_op.cpp
pderefmat_op.cpp
geom.cpp
gridfunc.cpp
hybridization.cpp
-41
View File
@@ -1275,22 +1275,6 @@ void BilinearForm::Update(FiniteElementSpace *nfes)
height = width = fes->GetVSize();
if (ext) { ext->Update(); }
for (int k = 0; k < domain_integs.Size(); ++k)
{
domain_integs[k]->Update();
}
for (int k = 0; k < boundary_integs.Size(); ++k)
{
boundary_integs[k]->Update();
}
for (int k = 0; k < interior_face_integs.Size(); ++k)
{
interior_face_integs[k]->Update();
}
for (int k = 0; k < boundary_integs.Size(); ++k)
{
boundary_face_integs[k]->Update();
}
}
void BilinearForm::SetDiagonalPolicy(DiagonalPolicy policy)
@@ -2353,31 +2337,6 @@ void MixedBilinearForm::Update()
height = test_fes->GetVSize();
width = trial_fes->GetVSize();
if (ext) { ext->Update(); }
for (int k = 0; k < domain_integs.Size(); ++k)
{
domain_integs[k]->Update();
}
for (int k = 0; k < boundary_integs.Size(); ++k)
{
boundary_integs[k]->Update();
}
for (int k = 0; k < interior_face_integs.Size(); ++k)
{
interior_face_integs[k]->Update();
}
for (int k = 0; k < boundary_integs.Size(); ++k)
{
boundary_face_integs[k]->Update();
}
for (int k = 0; k < trace_face_integs.Size(); ++k)
{
trace_face_integs[k]->Update();
}
for (int k = 0; k < boundary_trace_face_integs.Size(); ++k)
{
boundary_trace_face_integs[k]->Update();
}
}
MixedBilinearForm::~MixedBilinearForm()
+70 -26
View File
@@ -255,6 +255,8 @@ PABilinearFormExtension::PABilinearFormExtension(BilinearForm *form)
void PABilinearFormExtension::SetupRestrictionOperators(const L2FaceValues m)
{
MFEM_PERF_FUNCTION;
if ( Device::Allows(Backend::CEED_MASK) ) { return; }
ElementDofOrdering ordering = GetEVectorOrdering(*a->FESpace());
elem_restrict = trial_fes->GetElementRestriction(ordering);
@@ -266,7 +268,11 @@ void PABilinearFormExtension::SetupRestrictionOperators(const L2FaceValues m)
// Gather the attributes on the host from all the elements
const Mesh &mesh = *trial_fes->GetMesh();
elem_attributes = &mesh.GetElementAttributes();
elem_attributes.SetSize(mesh.GetNE());
for (int i = 0; i < mesh.GetNE(); ++i)
{
elem_attributes[i] = mesh.GetAttribute(i);
}
}
// Construct face restriction operators only if the bilinear form has
@@ -325,12 +331,52 @@ void PABilinearFormExtension::SetupRestrictionOperators(const L2FaceValues m)
bdr_face_dYdn.SetSize(bdr_face_restrict_lex->Height());
}
bdr_face_attributes = &trial_fes->GetMesh()->GetBdrFaceAttributes();
const Mesh &mesh = *trial_fes->GetMesh();
// See LinearFormExtension::Update for explanation of f_to_be logic.
std::unordered_map<int,int> f_to_be;
for (int i = 0; i < mesh.GetNBE(); ++i)
{
const int f = mesh.GetBdrElementFaceIndex(i);
f_to_be[f] = i;
}
const int nf_bdr = trial_fes->GetNFbyType(FaceType::Boundary);
bdr_attributes.SetSize(nf_bdr);
int f_ind = 0;
int missing_bdr_elems = 0;
for (int f = 0; f < mesh.GetNumFaces(); ++f)
{
if (!mesh.GetFaceInformation(f).IsOfFaceType(FaceType::Boundary))
{
continue;
}
int attribute = 1; // default value
if (f_to_be.find(f) != f_to_be.end())
{
const int be = f_to_be[f];
attribute = mesh.GetBdrAttribute(be);
}
else
{
// If a boundary face does not correspond to the a boundary element,
// we assign it the default attribute of 1. We also generate a
// warning at runtime with the number of such missing elements.
++missing_bdr_elems;
}
bdr_attributes[f_ind] = attribute;
++f_ind;
}
if (missing_bdr_elems)
{
MFEM_WARNING("Missing " << missing_bdr_elems << " boundary elements "
"for boundary faces.");
}
}
}
void PABilinearFormExtension::Assemble()
{
MFEM_PERF_FUNCTION;
SetupRestrictionOperators(L2FaceValues::DoubleValued);
Array<BilinearFormIntegrator*> &integrators = *a->GetDBFI();
@@ -387,7 +433,7 @@ void PABilinearFormExtension::AssembleDiagonal(Vector &y) const
mfem::forall(ne, [=] MFEM_HOST_DEVICE (int e)
{
const int attr = d_attr[e];
if (attr <= 0 || d_m[attr - 1] == 0)
if (d_m[attr - 1] == 0)
{
for (int i = 0; i < nd; ++i)
{
@@ -408,7 +454,7 @@ void PABilinearFormExtension::AssembleDiagonal(Vector &y) const
for (int i = 0; i < iSz; ++i)
{
assemble_diagonal_with_markers(*integrators[i], elem_markers[i],
*elem_attributes, localY);
elem_attributes, localY);
}
const ElementRestriction* H1elem_restrict =
dynamic_cast<const ElementRestriction*>(elem_restrict);
@@ -434,7 +480,7 @@ void PABilinearFormExtension::AssembleDiagonal(Vector &y) const
for (int i = 0; i < iSz; ++i)
{
assemble_diagonal_with_markers(*integrators[i], elem_markers[i],
*elem_attributes, y);
elem_attributes, y);
}
}
@@ -447,7 +493,7 @@ void PABilinearFormExtension::AssembleDiagonal(Vector &y) const
for (int i = 0; i < n_bdr_integs; ++i)
{
assemble_diagonal_with_markers(*bdr_integs[i], bdr_markers[i],
*bdr_face_attributes, bdr_face_Y);
bdr_attributes, bdr_face_Y);
}
bdr_face_restrict_lex->AddAbsMultTranspose(bdr_face_Y, y);
}
@@ -487,6 +533,8 @@ void PABilinearFormExtension::FormLinearSystem(const Array<int> &ess_tdof_list,
void PABilinearFormExtension::MultInternal(const Vector &x, Vector &y,
const bool useAbs) const
{
MFEM_PERF_FUNCTION;
Array<BilinearFormIntegrator*> &integrators = *a->GetDBFI();
const int iSz = integrators.Size();
@@ -546,7 +594,7 @@ void PABilinearFormExtension::MultInternal(const Vector &x, Vector &y,
for (int i = 0; i < iSz; ++i)
{
AddMultWithMarkers(*integrators[i], localX, elem_markers[i],
*elem_attributes, false, localY, useAbs);
elem_attributes, false, localY, useAbs);
}
if (H1elem_restrict && useAbs)
{
@@ -648,8 +696,8 @@ void PABilinearFormExtension::MultInternal(const Vector &x, Vector &y,
}
for (int i = 0; i < n_bdr_integs; ++i)
{
AddMultWithMarkers(*bdr_integs[i], bdr_face_X, bdr_markers[i],
*bdr_face_attributes, false, bdr_face_Y);
AddMultWithMarkers(*bdr_integs[i], bdr_face_X, bdr_markers[i], bdr_attributes,
false, bdr_face_Y);
}
for (int i = 0; i < n_bdr_face_integs; ++i)
{
@@ -657,14 +705,12 @@ void PABilinearFormExtension::MultInternal(const Vector &x, Vector &y,
{
AddMultNormalDerivativesWithMarkers(
*bdr_face_integs[i], bdr_face_X, bdr_face_dXdn,
bdr_face_markers[i], *bdr_face_attributes, bdr_face_Y,
bdr_face_dYdn);
bdr_face_markers[i], bdr_attributes, bdr_face_Y, bdr_face_dYdn);
}
else
{
AddMultWithMarkers(*bdr_face_integs[i], bdr_face_X,
bdr_face_markers[i], *bdr_face_attributes, false,
bdr_face_Y);
AddMultWithMarkers(*bdr_face_integs[i], bdr_face_X, bdr_face_markers[i],
bdr_attributes, false, bdr_face_Y);
}
}
bdr_face_restrict_lex->AddMultTransposeInPlace(bdr_face_Y, y);
@@ -687,7 +733,7 @@ void PABilinearFormExtension::MultTranspose(const Vector &x, Vector &y) const
localY = 0.0;
for (int i = 0; i < iSz; ++i)
{
AddMultWithMarkers(*integrators[i], localX, elem_markers[i], *elem_attributes,
AddMultWithMarkers(*integrators[i], localX, elem_markers[i], elem_attributes,
true, localY);
}
elem_restrict->MultTranspose(localY, y);
@@ -734,14 +780,13 @@ void PABilinearFormExtension::MultTranspose(const Vector &x, Vector &y) const
bdr_face_Y = 0.0;
for (int i = 0; i < n_bdr_integs; ++i)
{
AddMultWithMarkers(*bdr_integs[i], bdr_face_X, bdr_markers[i],
*bdr_face_attributes, true, bdr_face_Y);
AddMultWithMarkers(*bdr_integs[i], bdr_face_X, bdr_markers[i], bdr_attributes,
true, bdr_face_Y);
}
for (int i = 0; i < n_bdr_face_integs; ++i)
{
AddMultWithMarkers(*bdr_face_integs[i], bdr_face_X,
bdr_face_markers[i], *bdr_face_attributes, true,
bdr_face_Y);
AddMultWithMarkers(*bdr_face_integs[i], bdr_face_X, bdr_face_markers[i],
bdr_attributes, true, bdr_face_Y);
}
bdr_face_restrict_lex->AddMultTransposeInPlace(bdr_face_Y, y);
}
@@ -765,7 +810,7 @@ static void AddWithMarkers_(
mfem::forall(ne, [=] MFEM_HOST_DEVICE (int e)
{
const int attr = d_attr[e];
if (attr <= 0 || d_m[attr - 1] == 0) { return; }
if (d_m[attr - 1] == 0) { return; }
for (int i = 0; i < nd; ++i)
{
d_y(i, e) += d_x(i, e);
@@ -881,8 +926,7 @@ void EABilinearFormExtension::Assemble()
{
const int i = idx % sz;
const int e = idx / sz;
const real_t val =
d_a[e] > 0 ? (d_m[d_a[e] - 1] ? d_ea_1(i, e) : 0) : 0;
const real_t val = d_m[d_a[e] - 1] ? d_ea_1(i, e) : 0.0;
if (add)
{
d_ea_2(i, e) += val;
@@ -915,7 +959,7 @@ void EABilinearFormExtension::Assemble()
ea_data_tmp.SetSize(ea_data.Size());
integrators[i]->AssembleEA(*a->FESpace(), ea_data_tmp, false);
add_with_markers(ea_data_tmp, ea_data, ne, *markers,
*elem_attributes, add);
elem_attributes, add);
}
}
}
@@ -944,7 +988,7 @@ void EABilinearFormExtension::Assemble()
ea_data_tmp.SetSize(ea_data_bdr.Size());
bdr_integs[i]->AssembleEABoundary(*a->FESpace(), ea_data_tmp, add);
add_with_markers(ea_data_tmp, ea_data_bdr, nf_bdr, *markers,
*bdr_face_attributes, add);
bdr_attributes, add);
}
}
}
@@ -993,7 +1037,7 @@ void EABilinearFormExtension::Assemble()
ea_data_tmp,
add);
add_with_markers(ea_data_tmp, ea_data_bdr, nf_bdr, *markers,
*bdr_face_attributes, add);
bdr_attributes, add);
}
}
}
+1 -2
View File
@@ -69,8 +69,7 @@ class PABilinearFormExtension : public BilinearFormExtension
protected:
const FiniteElementSpace *trial_fes, *test_fes; // Not owned
/// Attributes of all mesh elements.
const Array<int> *elem_attributes; // Not owned
const Array<int> *bdr_face_attributes; // Not owned
Array<int> elem_attributes, bdr_attributes;
mutable Vector tmp_evec; // Work array
mutable Vector localX, localY;
mutable Vector int_face_X, int_face_Y;
-11
View File
@@ -21,11 +21,6 @@ using namespace std;
namespace mfem
{
void BilinearFormIntegrator::Update()
{
// default no-op
}
void BilinearFormIntegrator::AssemblePA(const FiniteElementSpace&)
{
MFEM_ABORT("BilinearFormIntegrator::AssemblePA(fes)\n"
@@ -3465,12 +3460,6 @@ real_t ElasticityIntegrator::ComputeFluxEnergy(const FiniteElement &fluxelem,
return energy;
}
void DGTraceIntegrator::Update()
{
qspace[0].reset();
qspace[1].reset();
}
void DGTraceIntegrator::AssembleFaceMatrix(const FiniteElement &el1,
const FiniteElement &el2,
FaceElementTransformations &Trans,
+10 -19
View File
@@ -23,8 +23,6 @@
namespace mfem
{
class QuadratureSpace;
class FaceQuadratureSpace;
/// Abstract base class BilinearFormIntegrator
class BilinearFormIntegrator : public NonlinearFormIntegrator
@@ -46,10 +44,6 @@ public:
// make sense for the action of the nonlinear operator (but they all make
// sense for its Jacobian).
/// Signal this integrator that something about either the trial or test space has changed.
virtual void Update();
/// Method defining partial assembly.
/** The result of the partial assembly is stored internally so that it can be
used later in the methods AddMultPA() and AddMultTransposePA(). */
@@ -818,7 +812,7 @@ protected:
const FiniteElement & test_fe) const
{
return (trial_fe.GetDim() == 1 && test_fe.GetDim() == 1 &&
trial_fe.GetDerivType() == mfem::FiniteElement::GRAD &&
trial_fe.GetDerivType() == mfem::FiniteElement::GRAD &&
test_fe.GetRangeType() == mfem::FiniteElement::SCALAR );
}
@@ -890,7 +884,7 @@ protected:
const FiniteElement & trial_fe,
const FiniteElement & test_fe) const
{
return (trial_fe.GetDerivType() == mfem::FiniteElement::DIV &&
return (trial_fe.GetDerivType() == mfem::FiniteElement::DIV &&
test_fe.GetRangeType() == mfem::FiniteElement::SCALAR );
}
@@ -925,7 +919,7 @@ protected:
const FiniteElement & trial_fe,
const FiniteElement & test_fe) const
{
return (trial_fe.GetDerivType() == mfem::FiniteElement::DIV &&
return (trial_fe.GetDerivType() == mfem::FiniteElement::DIV &&
test_fe.GetRangeType() == mfem::FiniteElement::VECTOR );
}
@@ -1606,7 +1600,7 @@ public:
{
return (trial_fe.GetCurlDim() == 3 && test_fe.GetRangeDim() == 3 &&
trial_fe.GetRangeType() == mfem::FiniteElement::VECTOR &&
trial_fe.GetDerivType() == mfem::FiniteElement::CURL &&
trial_fe.GetDerivType() == mfem::FiniteElement::CURL &&
test_fe.GetRangeType() == mfem::FiniteElement::VECTOR );
}
@@ -1641,7 +1635,7 @@ public:
{
return (trial_fe.GetDim() == 2 && test_fe.GetDim() == 2 &&
trial_fe.GetRangeType() == mfem::FiniteElement::VECTOR &&
trial_fe.GetDerivType() == mfem::FiniteElement::CURL &&
trial_fe.GetDerivType() == mfem::FiniteElement::CURL &&
test_fe.GetRangeType() == mfem::FiniteElement::VECTOR );
}
@@ -1675,7 +1669,7 @@ public:
{
return (trial_fe.GetDim() == 2 && test_fe.GetDim() == 2 &&
trial_fe.GetRangeType() == mfem::FiniteElement::SCALAR &&
trial_fe.GetDerivType() == mfem::FiniteElement::GRAD &&
trial_fe.GetDerivType() == mfem::FiniteElement::GRAD &&
test_fe.GetRangeType() == mfem::FiniteElement::SCALAR );
}
@@ -1766,7 +1760,7 @@ public:
const FiniteElement & test_fe) const
{
return (trial_fe.GetRangeType() == mfem::FiniteElement::SCALAR &&
trial_fe.GetDerivType() == mfem::FiniteElement::GRAD &&
trial_fe.GetDerivType() == mfem::FiniteElement::GRAD &&
test_fe.GetRangeType() == mfem::FiniteElement::SCALAR );
}
@@ -1799,7 +1793,7 @@ public:
const FiniteElement & test_fe) const
{
return (trial_fe.GetRangeType() == mfem::FiniteElement::SCALAR &&
trial_fe.GetDerivType() == mfem::FiniteElement::GRAD &&
trial_fe.GetDerivType() == mfem::FiniteElement::GRAD &&
test_fe.GetRangeType() == mfem::FiniteElement::VECTOR &&
test_fe.GetDerivType() == mfem::FiniteElement::DIV );
}
@@ -1838,7 +1832,7 @@ public:
const FiniteElement & test_fe) const
{
return (trial_fe.GetRangeType() == mfem::FiniteElement::VECTOR &&
trial_fe.GetDerivType() == mfem::FiniteElement::DIV &&
trial_fe.GetDerivType() == mfem::FiniteElement::DIV &&
test_fe.GetRangeType() == mfem::FiniteElement::SCALAR &&
test_fe.GetDerivType() == mfem::FiniteElement::GRAD
);
@@ -1979,7 +1973,7 @@ protected:
const FiniteElement & test_fe) const override
{
return (trial_fe.GetCurlDim() == 3 && test_fe.GetRangeDim() == 3 &&
trial_fe.GetDerivType() == mfem::FiniteElement::CURL &&
trial_fe.GetDerivType() == mfem::FiniteElement::CURL &&
test_fe.GetRangeType() == mfem::FiniteElement::VECTOR );
}
@@ -3317,7 +3311,6 @@ protected:
VectorCoefficient *u;
real_t alpha, beta;
// PA extension
std::unique_ptr<FaceQuadratureSpace> qspace[2];
Vector pa_data;
const DofToQuad *maps; ///< Not owned
const FaceGeometricFactors *geom; ///< Not owned
@@ -3340,8 +3333,6 @@ public:
real_t a, real_t b)
{ rho = &rho_; u = &u_; alpha = a; beta = b; }
void Update() override;
using BilinearFormIntegrator::AssembleFaceMatrix;
void AssembleFaceMatrix(const FiniteElement &el1,
const FiniteElement &el2,
+4 -4
View File
@@ -912,7 +912,7 @@ ConduitDataCollection::GridFunctionToBlueprintField(mfem::GridFunction *gf,
if (vdim == 1) // scalar case
{
n_field["values"].set_external(const_cast<real_t *>(gf->HostRead()),
n_field["values"].set_external(gf->GetData(),
ndofs);
}
else // vector case
@@ -925,18 +925,18 @@ ConduitDataCollection::GridFunctionToBlueprintField(mfem::GridFunction *gf,
int vdim_stride = (ordering == Ordering::byNODES ? ndofs : 1);
index_t offset = 0;
index_t stride = sizeof(real_t) * entry_stride;
index_t stride = sizeof(double) * entry_stride;
for (int d = 0; d < vdim; d++)
{
std::ostringstream oss;
oss << "v" << d;
std::string comp_name = oss.str();
n_field["values"][comp_name].set_external(const_cast<real_t *>(gf->HostRead()),
n_field["values"][comp_name].set_external(gf->GetData(),
ndofs,
offset,
stride);
offset += sizeof(real_t) * vdim_stride;
offset += sizeof(double) * vdim_stride;
}
}
-266
View File
@@ -1,266 +0,0 @@
// Copyright (c) 2010-2025, Lawrence Livermore National Security, LLC. Produced
// at the Lawrence Livermore National Laboratory. All Rights reserved. See files
// LICENSE and NOTICE for details. LLNL-CODE-806117.
//
// This file is part of the MFEM library. For more information and source code
// availability visit https://mfem.org.
//
// MFEM is free software; you can redistribute it and/or modify it under the
// terms of the BSD-3 license. We welcome feedback and contributions, see file
// CONTRIBUTING.md for details.
#include "derefmat_op.hpp"
#include "fes_kernels.hpp"
/// \cond DO_NOT_DOCUMENT
namespace mfem
{
namespace internal
{
template <Ordering::Type Order, bool Atomic>
static void DerefMultKernelImpl(const DerefineMatrixOp &op, const Vector &x,
Vector &y)
{
DerefineMatrixOpMultFunctor<Order, Atomic> func;
func.xptr = x.Read();
y.UseDevice();
y = 0.;
func.yptr = y.ReadWrite();
func.bsptr = op.block_storage.Read();
func.boptr = op.block_offsets.Read();
func.brptr = op.block_row_idcs_offsets.Read();
func.bcptr = op.block_col_idcs_offsets.Read();
func.rptr = op.row_idcs.Read();
func.cptr = op.col_idcs.Read();
func.vdims = op.fespace->GetVDim();
func.nblocks = op.block_offsets.Size();
func.width = op.Width() / func.vdims;
func.height = op.Height() / func.vdims;
func.Run(op.max_rows);
}
} // namespace internal
DerefineMatrixOp::DerefineMatrixOp(FiniteElementSpace &fespace_, int old_ndofs,
const Table *old_elem_dof,
const Table *old_elem_fos)
: Operator(fespace_.GetVSize(), old_ndofs * fespace_.GetVDim()),
fespace(&fespace_)
{
static Kernels kernels;
constexpr int max_team_size = 256;
/// TODO: Implement DofTransformation support
MFEM_VERIFY(fespace->Nonconforming(),
"Not implemented for conforming meshes.");
MFEM_VERIFY(old_ndofs, "Missing previous (finer) space.");
MFEM_VERIFY(fespace->GetNDofs() <= old_ndofs,
"Previous space is not finer.");
const CoarseFineTransformations &dtrans =
fespace->GetMesh()->ncmesh->GetDerefinementTransforms();
MFEM_ASSERT(dtrans.embeddings.Size() == old_elem_dof->Size(), "");
const bool is_dg = fespace->FEColl()->GetContType()
== FiniteElementCollection::DISCONTINUOUS;
DenseMatrix localRVO; // for variable-order only
DenseTensor localR[Geometry::NumGeom];
int total_rows = 0;
int total_cols = 0;
block_offsets.SetSize(dtrans.embeddings.Size());
block_offsets.HostWrite();
if (fespace->IsVariableOrder())
{
// TODO: any potential for some compression here?
// determine storage size and offsets
block_offsets[0] = 0;
int total_size = 0;
for (int k = 0; k < dtrans.embeddings.Size(); ++k)
{
const Embedding &emb = dtrans.embeddings[k];
const FiniteElement *fe = fespace->GetFE(emb.parent);
const int ldof = fe->GetDof();
if (k + 1 < dtrans.embeddings.Size())
{
block_offsets[k + 1] = block_offsets[k] + ldof * ldof;
}
total_rows += ldof;
total_cols += ldof;
total_size += ldof * ldof;
}
block_storage.SetSize(total_size);
}
else
{
// compression scheme:
// block_offsets is the start of each block, potentially repeated
// only need to store localR for used shapes
Mesh::GeometryList elem_geoms(*fespace->GetMesh());
int geom_offsets[Geometry::NumGeom];
{
int size = 0;
for (int i = 0; i < elem_geoms.Size(); ++i)
{
fespace->GetLocalDerefinementMatrices(elem_geoms[i],
localR[elem_geoms[i]]);
geom_offsets[elem_geoms[i]] = size;
size += localR[elem_geoms[i]].TotalSize();
}
block_storage.SetSize(size);
// copy blocks into block_storage
auto bs_ptr = block_storage.HostWrite();
for (int i = 0; i < elem_geoms.Size(); ++i)
{
std::copy(localR[elem_geoms[i]].Data(),
localR[elem_geoms[i]].Data()
+ localR[elem_geoms[i]].TotalSize(),
bs_ptr);
bs_ptr += localR[elem_geoms[i]].TotalSize();
}
}
for (int k = 0; k < dtrans.embeddings.Size(); ++k)
{
const Embedding &emb = dtrans.embeddings[k];
Geometry::Type geom =
fespace->GetMesh()->GetElementBaseGeometry(emb.parent);
auto size = localR[geom].SizeI() * localR[geom].SizeJ();
total_rows += localR[geom].SizeI();
total_cols += localR[geom].SizeJ();
// set block offsets and sizes
block_offsets[k] = geom_offsets[geom] + size * emb.matrix;
}
}
row_idcs.SetSize(total_rows);
row_idcs.HostWrite();
col_idcs.SetSize(total_cols);
col_idcs.HostWrite();
block_row_idcs_offsets.SetSize(dtrans.embeddings.Size() + 1);
block_row_idcs_offsets.HostWrite();
block_col_idcs_offsets.SetSize(dtrans.embeddings.Size() + 1);
block_col_idcs_offsets.HostWrite();
block_row_idcs_offsets[0] = 0;
block_col_idcs_offsets[0] = 0;
// compute index information
Array<int> dofs, old_dofs;
max_rows = 1;
{
Array<int> mark(fespace->GetNDofs());
mark = 0;
auto bs_ptr = block_storage.HostWrite();
int ridx = 0;
int cidx = 0;
int num_marked = 0;
for (int k = 0; k < dtrans.embeddings.Size(); k++)
{
const Embedding &emb = dtrans.embeddings[k];
Geometry::Type geom =
fespace->GetMesh()->GetElementBaseGeometry(emb.parent);
if (fespace->IsVariableOrder())
{
const FiniteElement *fe = fespace->GetFE(emb.parent);
const DenseTensor &pmats = dtrans.point_matrices[geom];
const int ldof = fe->GetDof();
IsoparametricTransformation isotr;
isotr.SetIdentityTransformation(geom);
localRVO.SetSize(ldof, ldof);
isotr.SetPointMat(pmats(emb.matrix));
// Local restriction is size ldofxldof assuming that the parent
// and child are of same polynomial order.
fe->GetLocalRestriction(isotr, localRVO);
// copy block
auto size = localRVO.Height() * localRVO.Width();
std::copy(localRVO.Data(), localRVO.Data() + size, bs_ptr);
bs_ptr += size;
}
DenseMatrix &lR =
fespace->IsVariableOrder() ? localRVO : localR[geom](emb.matrix);
block_row_idcs_offsets[k + 1] =
block_row_idcs_offsets[k] + lR.Height();
block_col_idcs_offsets[k + 1] = block_col_idcs_offsets[k] + lR.Width();
max_rows = std::max(lR.Height(), max_rows);
// index information
fespace->elem_dof->GetRow(emb.parent, dofs);
old_elem_dof->GetRow(k, old_dofs);
MFEM_VERIFY(old_dofs.Size() == dofs.Size(),
"Parent and child must have same #dofs.");
for (int i = 0; i < lR.Height(); ++i, ++ridx)
{
if (!std::isfinite(lR(i, 0)))
{
row_idcs[ridx] = INT_MAX;
continue;
}
int r = dofs[i];
int m = (r >= 0) ? r : (-1 - r);
if (is_dg || !mark[m])
{
row_idcs[ridx] = r;
mark[m] = 1;
++num_marked;
}
else
{
row_idcs[ridx] = INT_MAX;
}
}
for (int i = 0; i < lR.Width(); ++i, ++cidx)
{
col_idcs[cidx] = old_dofs[i];
}
}
if (!is_dg && !fespace->IsVariableOrder())
{
MFEM_VERIFY(num_marked * fespace->GetVDim() == Height(),
"internal error: not all rows were set.");
}
}
// if not using GPU, set max_rows/max_cols to zero
if (Device::Allows(Backend::DEVICE_MASK))
{
max_rows = std::min(max_rows, max_team_size);
}
else
{
max_rows = 1;
}
}
void DerefineMatrixOp::Mult(const Vector &x, Vector &y) const
{
const bool is_dg = fespace->FEColl()->GetContType()
== FiniteElementCollection::DISCONTINUOUS;
// DG needs atomic summation
MultKernel::Run(fespace->GetOrdering(), is_dg, *this, x, y);
}
DerefineMatrixOp::Kernels::Kernels()
{
MultKernel::Specialization<Ordering::byNODES, false>::Add();
MultKernel::Specialization<Ordering::byVDIM, false>::Add();
MultKernel::Specialization<Ordering::byNODES, true>::Add();
MultKernel::Specialization<Ordering::byVDIM, true>::Add();
}
template <Ordering::Type Order, bool Atomic>
DerefineMatrixOp::MultKernelType DerefineMatrixOp::MultKernel::Kernel()
{
return internal::DerefMultKernelImpl<Order, Atomic>;
}
DerefineMatrixOp::MultKernelType
DerefineMatrixOp::MultKernel::Fallback(Ordering::Type, bool)
{
MFEM_ABORT("invalid MultKernel parameters");
}
} // namespace mfem
/// \endcond DO_NOT_DOCUMENT
-65
View File
@@ -1,65 +0,0 @@
// Copyright (c) 2010-2025, Lawrence Livermore National Security, LLC. Produced
// at the Lawrence Livermore National Laboratory. All Rights reserved. See files
// LICENSE and NOTICE for details. LLNL-CODE-806117.
//
// This file is part of the MFEM library. For more information and source code
// availability visit https://mfem.org.
//
// MFEM is free software; you can redistribute it and/or modify it under the
// terms of the BSD-3 license. We welcome feedback and contributions, see file
// CONTRIBUTING.md for details.
#ifndef MFEM_DEREFMAT_OP
#define MFEM_DEREFMAT_OP
#include "fespace.hpp"
#include "kernel_dispatch.hpp"
namespace mfem
{
/// \cond DO_NOT_DOCUMENT
struct DerefineMatrixOp : public Operator
{
FiniteElementSpace *fespace;
/// offsets into block_storage
Array<int> block_offsets;
/// offsets into row_idcs
Array<int> block_row_idcs_offsets;
/// offsets into col_idcs
Array<int> block_col_idcs_offsets;
/// mapping for row dofs, INT_MAX indicates the block row should be ignored.
/// negative means the row data should be negated.
Array<int> row_idcs;
/// mapping for col dofs, negative means the col data should be negated.
Array<int> col_idcs;
/// dense block matrices which can be reused to construct the full matrix
/// operation. These are stored contiguously and blocks have no restrictions
/// on shape (can be rectangle and differ from block to block).
Vector block_storage;
/// maximum height of any block in block_storage for GPU
/// parallelization, or 1 for CPU runs.
int max_rows;
using MultKernelType = void (*)(const DerefineMatrixOp &, const Vector &,
Vector &);
/// template args: ordering, atomic
MFEM_REGISTER_KERNELS(MultKernel, MultKernelType, (Ordering::Type, bool));
struct Kernels
{
Kernels();
};
void Mult(const Vector &x, Vector &y) const;
DerefineMatrixOp(FiniteElementSpace &fespace_, int old_ndofs,
const Table *old_elem_dof, const Table *old_elem_fos);
};
/// \endcond DO_NOT_DOCUMENT
} // namespace mfem
#endif
-1
View File
@@ -241,7 +241,6 @@ public:
{
MFEM_ASSERT(!action_callbacks.empty(), "no integrators have been set");
prolongation(solutions, solutions_t, solutions_l);
residual_l = 0.0;
for (auto &action : action_callbacks)
{
action(solutions_l, parameters_l, residual_l);
+6 -6
View File
@@ -327,8 +327,8 @@ void print_mpi_sync(const std::string& msg)
// First gather string lengths
size_t msg_len = msg.length();
std::vector<size_t> lengths(nranks);
MPI_Gather(&msg_len, 1, MPITypeMap<size_t>::mpi_type,
lengths.data(), 1, MPITypeMap<size_t>::mpi_type,
MPI_Gather(&msg_len, 1, MPI_INT,
lengths.data(), 1, MPI_INT,
0, MPI_COMM_WORLD);
if (myrank == 0)
@@ -568,7 +568,7 @@ struct ThreadBlocks
int z = 1;
};
#if defined(MFEM_USE_CUDA_OR_HIP)
#if (defined(MFEM_USE_CUDA) || defined(MFEM_USE_HIP))
template <typename func_t>
__global__ void forall_kernel_shmem(func_t f, int n)
{
@@ -591,7 +591,7 @@ void forall(func_t f,
if (Device::Allows(Backend::CUDA_MASK) ||
Device::Allows(Backend::HIP_MASK))
{
#if defined(MFEM_USE_CUDA_OR_HIP)
#if (defined(MFEM_USE_CUDA) || defined(MFEM_USE_HIP))
// int gridsize = (N + Z - 1) / Z;
int num_bytes = num_shmem * sizeof(decltype(shmem));
dim3 block_size(blocks.x, blocks.y, blocks.z);
@@ -987,7 +987,7 @@ get_restriction_transpose(
{
auto RT = [=](const Vector &v_e, Vector &v_l)
{
v_l += v_e;
v_l = v_e;
};
return std::make_tuple(RT, 1);
}
@@ -996,7 +996,7 @@ get_restriction_transpose(
const Operator *R = get_restriction<entity_t>(f, o);
std::function<void(const Vector&, Vector&)> RT = [=](const Vector &x, Vector &y)
{
R->AddMultTranspose(x, y);
R->MultTranspose(x, y);
};
return std::make_tuple(RT, R->Height());
}
-249
View File
@@ -1,249 +0,0 @@
// Copyright (c) 2010-2025, Lawrence Livermore National Security, LLC. Produced
// at the Lawrence Livermore National Laboratory. All Rights reserved. See files
// LICENSE and NOTICE for details. LLNL-CODE-806117.
//
// This file is part of the MFEM library. For more information and source code
// availability visit https://mfem.org.
//
// MFEM is free software; you can redistribute it and/or modify it under the
// terms of the BSD-3 license. We welcome feedback and contributions, see file
// CONTRIBUTING.md for details.
#ifndef MFEM_FES_KERNELS_HPP
#define MFEM_FES_KERNELS_HPP
#include "../general/forall.hpp"
#include <climits>
namespace mfem
{
/// \cond DO_NOT_DOCUMENT
namespace internal
{
///
/// Implements matrix-vector multiply $y = A x$ for a sparse matrix composed of
/// a sum of smaller dense blocks. There is additional permutation/sign
/// information associated with each block. The base class only implements
/// helper routines such as computing block widths, index into x, index into y,
/// and column in A given sub-block information.
/// @sa DerefineMatrixOpMultFunctor
///
/// @tparam Order vdim ordering for x and y. Note that for Diag = false this is
/// ignored for x as x has a special interleaved order.
/// @tparam Base used for the curious recurring template pattern (CRTP) so the
/// base class can access child class fields without virtual functions
/// @tparam Diag true if this corresponds to the diagonal block (coarse element
/// and fine element are on our rank), false otherwise (coarse element is on our
/// rank, fine element is on a different rank).
///
template <Ordering::Type Order, class Base, bool Diag = true>
struct DerefineMatrixOpFunctorBase;
template <class Base>
struct DerefineMatrixOpFunctorBase<Ordering::byNODES, Base, true>
{
/// block column indices offsets
const int *bcptr;
/// column indices
const int *cptr;
int MFEM_HOST_DEVICE BlockWidth(int k) const
{
return bcptr[k + 1] - bcptr[k];
}
void MFEM_HOST_DEVICE Col(int j, int k, int &col, int &sign) const
{
col = cptr[bcptr[k] + j];
if (col < 0)
{
col = -1 - col;
sign = -sign;
}
}
int MFEM_HOST_DEVICE IndexX(int col, int vdim, int) const
{
return col + vdim * static_cast<const Base *>(this)->width;
}
int MFEM_HOST_DEVICE IndexY(int row, int vdim) const
{
return row + vdim * static_cast<const Base *>(this)->height;
}
};
template <class Base>
struct DerefineMatrixOpFunctorBase<Ordering::byVDIM, Base, true>
{
/// block column indices offsets
const int *bcptr;
/// column indices
const int *cptr;
int MFEM_HOST_DEVICE BlockWidth(int k) const
{
return bcptr[k + 1] - bcptr[k];
}
void MFEM_HOST_DEVICE Col(int j, int k, int &col, int &sign) const
{
col = cptr[bcptr[k] + j];
if (col < 0)
{
col = -1 - col;
sign = -sign;
}
}
int MFEM_HOST_DEVICE IndexX(int col, int vdim, int) const
{
return vdim + col * static_cast<const Base *>(this)->vdims;
}
int MFEM_HOST_DEVICE IndexY(int row, int vdim) const
{
return vdim + row * static_cast<const Base *>(this)->vdims;
}
};
template <class Base>
struct DerefineMatrixOpFunctorBase<Ordering::byNODES, Base, false>
{
/// receive segment offsets
const int *segptr;
/// receive segment index
const int *rsptr;
/// off-diagonal block column offsets
const int *coptr;
/// off-diagonal block widths
const int *bwptr;
int MFEM_HOST_DEVICE BlockWidth(int k) const { return bwptr[k]; }
void MFEM_HOST_DEVICE Col(int j, int k, int &col, int &sign) const
{
col = coptr[k] + j;
}
int MFEM_HOST_DEVICE IndexX(int col, int vdim, int k) const
{
int tmp = rsptr[k];
int segwidth = segptr[tmp + 1] - segptr[tmp];
return segptr[tmp] * static_cast<const Base *>(this)->vdims + col +
vdim * segwidth;
}
int MFEM_HOST_DEVICE IndexY(int row, int vdim) const
{
return row + vdim * static_cast<const Base *>(this)->height;
}
};
template <class Base>
struct DerefineMatrixOpFunctorBase<Ordering::byVDIM, Base, false>
{
/// receive segment offsets
const int *segptr;
/// receive segment index
const int *rsptr;
/// off-diagonal block column offsets
const int *coptr;
/// off-diagonal block widths
const int *bwptr;
int MFEM_HOST_DEVICE BlockWidth(int k) const { return bwptr[k]; }
void MFEM_HOST_DEVICE Col(int j, int k, int &col, int &sign) const
{
col = coptr[k] + j;
}
int MFEM_HOST_DEVICE IndexX(int col, int vdim, int k) const
{
int tmp = rsptr[k];
int segwidth = segptr[tmp + 1] - segptr[tmp];
return segptr[tmp] * static_cast<const Base *>(this)->vdims + col +
vdim * segwidth;
}
int MFEM_HOST_DEVICE IndexY(int row, int vdim) const
{
return vdim + row * static_cast<const Base *>(this)->vdims;
}
};
/// internally used to implement the derefinement operator Mult diagonal
/// block
template <Ordering::Type Order, bool Atomic, bool Diag = true>
struct DerefineMatrixOpMultFunctor
: public DerefineMatrixOpFunctorBase<
Order, DerefineMatrixOpMultFunctor<Order, Atomic, Diag>, Diag>
{
const real_t *xptr;
real_t *yptr;
/// block storage
const real_t *bsptr;
/// block offsets
const int *boptr;
/// block row index offsets
const int *brptr;
/// row indices
const int *rptr;
// number of blocks
int nblocks;
// number of components
int vdims;
/// overall operator height (for vdim = 1)
int height;
/// overall operator width (for vdim = 1)
int width;
void MFEM_HOST_DEVICE operator()(int kidx) const
{
int k = kidx % nblocks;
int vdim = kidx / nblocks;
int block_height = brptr[k + 1] - brptr[k];
int block_width = this->BlockWidth(k);
MFEM_FOREACH_THREAD(i, x, block_height)
{
int row = rptr[brptr[k] + i];
int rsign = 1;
if (row < 0)
{
row = -1 - row;
rsign = -1;
}
if (row < INT_MAX)
{
// row not marked as unused
real_t sum = 0;
for (int j = 0; j < block_width; ++j)
{
int col, sign = rsign;
this->Col(j, k, col, sign);
sum += sign * bsptr[boptr[k] + i + j * block_height] *
xptr[this->IndexX(col, vdim, k)];
}
#if defined(__CUDA_ARCH__) or defined(__HIP_DEVICE_COMPILE__)
if (Atomic)
{
atomicAdd(yptr + this->IndexY(row, vdim), sum);
}
else
#endif
{
yptr[this->IndexY(row, vdim)] += sum;
}
}
}
}
/// N is the max block row size (doesn't have to be a power of 2)
void Run(int N) const { forall_2D(nblocks * vdims, N, 1, *this); }
};
} // namespace internal
/// \endcond DO_NOT_DOCUMENT
} // namespace mfem
#endif
+6 -13
View File
@@ -17,9 +17,6 @@
#include "fem.hpp"
#include "ceed/interface/util.hpp"
#include "derefmat_op.hpp"
#include <algorithm>
#include <cmath>
#include <cstdarg>
@@ -27,9 +24,9 @@ using namespace std;
namespace mfem
{
template <>
void Ordering::DofsToVDofs<Ordering::byNODES>(int ndofs, int vdim,
Array<int> &dofs)
template <> void Ordering::
DofsToVDofs<Ordering::byNODES>(int ndofs, int vdim, Array<int> &dofs)
{
// static method
int size = dofs.Size();
@@ -43,9 +40,8 @@ void Ordering::DofsToVDofs<Ordering::byNODES>(int ndofs, int vdim,
}
}
template <>
void Ordering::DofsToVDofs<Ordering::byVDIM>(int ndofs, int vdim,
Array<int> &dofs)
template <> void Ordering::
DofsToVDofs<Ordering::byVDIM>(int ndofs, int vdim, Array<int> &dofs)
{
// static method
int size = dofs.Size();
@@ -59,6 +55,7 @@ void Ordering::DofsToVDofs<Ordering::byVDIM>(int ndofs, int vdim,
}
}
FiniteElementSpace::FiniteElementSpace()
: mesh(NULL), fec(NULL), vdim(0), ordering(Ordering::byNODES),
ndofs(0), nvdofs(0), nedofs(0), nfdofs(0), nbdofs(0),
@@ -4247,11 +4244,7 @@ void FiniteElementSpace::Update(bool want_transform)
case Mesh::DEREFINE:
{
BuildConformingInterpolation();
#if 0
Th.Reset(DerefinementMatrix(old_ndofs, old_elem_dof, old_elem_fos));
#else
Th.Reset(new DerefineMatrixOp(*this, old_ndofs, old_elem_dof, old_elem_fos));
#endif
if (IsVariableOrder())
{
if (cP && cR_hp)
+1 -2
View File
@@ -113,7 +113,7 @@ class QuadratureSpace;
class QuadratureInterpolator;
class FaceQuadratureInterpolator;
class PRefinementTransferOperator;
struct DerefineMatrixOp;
/** @brief Class FiniteElementSpace - responsible for providing FEM view of the
mesh, mainly managing the set of degrees of freedom.
@@ -246,7 +246,6 @@ class FiniteElementSpace
friend class PRefinementTransferOperator;
friend void Mesh::Swap(Mesh &, bool);
friend class LORBase;
friend struct DerefineMatrixOp;
protected:
/// The mesh that FE space lives on (not owned).
+4 -11
View File
@@ -139,6 +139,8 @@ void DGTraceIntegrator::SetupPA(const FiniteElementSpace &fes, FaceType type)
const MemoryType mt = (pa_mt == MemoryType::DEFAULT) ?
Device::GetDeviceMemoryType() : pa_mt;
nf = fes.GetNFbyType(type);
if (nf==0) { return; }
// Assumes tensor-product elements
Mesh *mesh = fes.GetMesh();
const FiniteElement &el = *fes.GetTypicalTraceElement();
@@ -146,17 +148,6 @@ void DGTraceIntegrator::SetupPA(const FiniteElementSpace &fes, FaceType type)
IntRule:
&GetRule(el.GetGeomType(), el.GetOrder(),
*mesh->GetTypicalElementTransformation());
if (!qspace[static_cast<int>(type)])
{
qspace[static_cast<int>(type)].reset(
new FaceQuadratureSpace(*mesh, *ir, type));
}
FaceQuadratureSpace& qs = *qspace[static_cast<int>(type)];
nf = qs.GetNumFaces();
if (nf==0) { return; }
const int symmDims = 4;
nq = ir->GetNPoints();
dim = mesh->Dimension();
@@ -168,6 +159,8 @@ void DGTraceIntegrator::SetupPA(const FiniteElementSpace &fes, FaceType type)
dofs1D = maps->ndof;
quad1D = maps->nqpt;
pa_data.SetSize(symmDims * nq * nf, Device::GetMemoryType());
FaceQuadratureSpace qs(*mesh, *ir, type);
CoefficientVector vel(*u, qs, CoefficientStorage::COMPRESSED);
CoefficientVector r(qs, CoefficientStorage::COMPRESSED);
+7 -10
View File
@@ -1212,22 +1212,18 @@ inline void SmemPADiffusionApply3D(const int NE,
} // namespace internal
namespace
{
using ApplyKernelType = DiffusionIntegrator::ApplyKernelType;
using DiagonalKernelType = DiffusionIntegrator::DiagonalKernelType;
}
template<int DIM, int T_D1D, int T_Q1D>
ApplyKernelType DiffusionIntegrator::ApplyPAKernels::Kernel()
DiffusionIntegrator::ApplyKernelType
DiffusionIntegrator::ApplyPAKernels::Kernel()
{
if (DIM == 2) { return internal::SmemPADiffusionApply2D<T_D1D,T_Q1D>; }
else if (DIM == 3) { return internal::SmemPADiffusionApply3D<T_D1D, T_Q1D>; }
else { MFEM_ABORT(""); }
}
inline
ApplyKernelType DiffusionIntegrator::ApplyPAKernels::Fallback(int DIM, int, int)
inline DiffusionIntegrator::ApplyKernelType
DiffusionIntegrator::ApplyPAKernels::Fallback(int DIM, int, int)
{
if (DIM == 2) { return internal::PADiffusionApply2D; }
else if (DIM == 3) { return internal::PADiffusionApply3D; }
@@ -1235,14 +1231,15 @@ ApplyKernelType DiffusionIntegrator::ApplyPAKernels::Fallback(int DIM, int, int)
}
template<int DIM, int D1D, int Q1D>
DiagonalKernelType DiffusionIntegrator::DiagonalPAKernels::Kernel()
DiffusionIntegrator::DiagonalKernelType
DiffusionIntegrator::DiagonalPAKernels::Kernel()
{
if (DIM == 2) { return internal::SmemPADiffusionDiagonal2D<D1D,Q1D>; }
else if (DIM == 3) { return internal::SmemPADiffusionDiagonal3D<D1D, Q1D>; }
else { MFEM_ABORT(""); }
}
inline DiagonalKernelType
inline DiffusionIntegrator::DiagonalKernelType
DiffusionIntegrator::DiagonalPAKernels::Fallback(int DIM, int, int)
{
if (DIM == 2) { return internal::PADiffusionDiagonal2D; }
+4
View File
@@ -39,6 +39,8 @@ void DiffusionIntegrator::AssembleDiagonalPA(Vector &diag)
// PA Diffusion Apply kernel
void DiffusionIntegrator::AddMultPA(const Vector &x, Vector &y) const
{
MFEM_PERF_FUNCTION;
if (DeviceCanUseCeed())
{
ceedOp->AddMult(x, y);
@@ -88,6 +90,8 @@ void DiffusionIntegrator::AddMultTransposePA(const Vector &x, Vector &y) const
void DiffusionIntegrator::AssemblePA(const FiniteElementSpace &fes)
{
MFEM_PERF_FUNCTION;
const MemoryType mt = (pa_mt == MemoryType::DEFAULT) ?
Device::GetDeviceMemoryType() : pa_mt;
// Assuming the same element type
+6 -11
View File
@@ -1383,14 +1383,9 @@ inline void EAMassAssemble3D(const int NE,
} // namespace internal
namespace
{
using ApplyKernelType = MassIntegrator::ApplyKernelType;
using DiagonalKernelType = MassIntegrator::DiagonalKernelType;
}
template<int DIM, int T_D1D, int T_Q1D>
ApplyKernelType MassIntegrator::ApplyPAKernels::Kernel()
MassIntegrator::ApplyKernelType MassIntegrator::ApplyPAKernels::Kernel()
{
if (DIM == 1) { return internal::PAMassApply1D; }
else if (DIM == 2) { return internal::SmemPAMassApply2D<T_D1D,T_Q1D>; }
@@ -1398,8 +1393,8 @@ ApplyKernelType MassIntegrator::ApplyPAKernels::Kernel()
else { MFEM_ABORT(""); }
}
inline ApplyKernelType MassIntegrator::ApplyPAKernels::Fallback(
int DIM, int, int)
inline MassIntegrator::ApplyKernelType
MassIntegrator::ApplyPAKernels::Fallback(int DIM, int, int)
{
if (DIM == 1) { return internal::PAMassApply1D; }
else if (DIM == 2) { return internal::PAMassApply2D; }
@@ -1408,7 +1403,7 @@ inline ApplyKernelType MassIntegrator::ApplyPAKernels::Fallback(
}
template<int DIM, int T_D1D, int T_Q1D>
DiagonalKernelType MassIntegrator::DiagonalPAKernels::Kernel()
MassIntegrator::DiagonalKernelType MassIntegrator::DiagonalPAKernels::Kernel()
{
if (DIM == 1) { return internal::PAMassAssembleDiagonal1D; }
else if (DIM == 2) { return internal::SmemPAMassAssembleDiagonal2D<T_D1D,T_Q1D>; }
@@ -1416,8 +1411,8 @@ DiagonalKernelType MassIntegrator::DiagonalPAKernels::Kernel()
else { MFEM_ABORT(""); }
}
inline DiagonalKernelType MassIntegrator::DiagonalPAKernels::Fallback(
int DIM, int, int)
inline MassIntegrator::DiagonalKernelType
MassIntegrator::DiagonalPAKernels::Fallback(int DIM, int, int)
{
if (DIM == 1) { return internal::PAMassAssembleDiagonal1D; }
else if (DIM == 2) { return internal::PAMassAssembleDiagonal2D; }
+60 -25
View File
@@ -23,6 +23,8 @@ namespace mfem
void MassIntegrator::AssemblePA(const FiniteElementSpace &fes)
{
MFEM_PERF_FUNCTION;
const MemoryType mt = (pa_mt == MemoryType::DEFAULT) ?
Device::GetDeviceMemoryType() : pa_mt;
@@ -59,23 +61,26 @@ void MassIntegrator::AssemblePA(const FiniteElementSpace &fes)
QuadratureSpace qs(*mesh, *ir);
CoefficientVector coeff(Q, qs, CoefficientStorage::COMPRESSED);
const int NE = ne;
const int Q1D = quad1D;
const int NQ = static_cast<int>(std::pow(Q1D, dim));
const bool const_c = coeff.Size() == 1;
const bool by_val = map_type == FiniteElement::VALUE;
const auto W = Reshape(ir->GetWeights().Read(), NQ);
const auto J = Reshape(geom->detJ.Read(), NQ, NE);
const auto C = const_c ? Reshape(coeff.Read(), 1, 1) :
Reshape(coeff.Read(), NQ,NE);
auto v = Reshape(pa_data.Write(), NQ, NE);
mfem::forall_2D(NE, NQ, 1, [=] MFEM_HOST_DEVICE (int e)
{
const int NE = ne;
const int NQ = nq;
const bool const_c = coeff.Size() == 1;
const bool by_val = map_type == FiniteElement::VALUE;
const auto W = Reshape(ir->GetWeights().Read(), NQ);
const auto J = Reshape(geom->detJ.Read(), NQ, NE);
const auto C =
const_c ? Reshape(coeff.Read(), 1, 1) : Reshape(coeff.Read(), NQ, NE);
auto v = Reshape(pa_data.Write(), NQ, NE);
mfem::forall(NQ, NE, [=] MFEM_HOST_DEVICE(int q, int e)
MFEM_FOREACH_THREAD(i, x, NQ)
{
const real_t detJ = J(q, e);
const real_t coeff = const_c ? C(0, 0) : C(q, e);
v(q, e) = W(q) * coeff * (by_val ? detJ : 1.0 / detJ);
});
}
const real_t detJ = J(i,e);
const real_t coeff = const_c ? C(0,0) : C(i,e);
v(i,e) = W(i) * coeff * (by_val ? detJ : 1.0/detJ);
}
});
}
void MassIntegrator::AssemblePABoundary(const FiniteElementSpace &fes)
@@ -106,22 +111,50 @@ void MassIntegrator::AssemblePABoundary(const FiniteElementSpace &fes)
CoefficientVector coeff(Q, qs, CoefficientStorage::COMPRESSED);
const int NE = ne;
const int NQ = nq;
const int Q1D = quad1D;
const bool const_c = coeff.Size() == 1;
const bool by_val = map_type == FiniteElement::VALUE;
if (dim==1)
{
const auto W = Reshape(ir->GetWeights().Read(), NQ);
const auto J = Reshape(face_geom->detJ.Read(), NQ, NE);
const auto C = const_c ? Reshape(coeff.Read(), 1, 1)
: Reshape(coeff.Read(), NQ, NE);
auto v = Reshape(pa_data.Write(), NQ, NE);
mfem::forall(NQ, NE, [=] MFEM_HOST_DEVICE(int q, int e)
const auto W = Reshape(ir->GetWeights().Read(), Q1D);
const auto J = Reshape(face_geom->detJ.Read(), Q1D, NE);
const auto C = const_c ? Reshape(coeff.Read(), 1, 1) :
Reshape(coeff.Read(), Q1D, NE);
auto v = Reshape(pa_data.Write(), Q1D, NE);
mfem::forall_2D(NE, Q1D, 1, [=] MFEM_HOST_DEVICE (int e)
{
const real_t detJ = J(q, e);
const real_t coeff = const_c ? C(0, 0) : C(q, e);
v(q, e) = W(q) * coeff * (by_val ? detJ : 1.0 / detJ);
MFEM_FOREACH_THREAD(qx,x,Q1D)
{
const real_t detJ = J(qx,e);
const real_t coeff = const_c ? C(0,0) : C(qx,e);
v(qx,e) = W(qx) * coeff * (by_val ? detJ : 1.0/detJ);
}
});
}
else if (dim==2)
{
const auto W = Reshape(ir->GetWeights().Read(), Q1D,Q1D);
const auto J = Reshape(face_geom->detJ.Read(), Q1D,Q1D,NE);
const auto C = const_c ? Reshape(coeff.Read(), 1,1,1) :
Reshape(coeff.Read(), Q1D,Q1D,NE);
auto v = Reshape(pa_data.Write(), Q1D,Q1D, NE);
mfem::forall_2D(NE, Q1D, Q1D, [=] MFEM_HOST_DEVICE (int e)
{
MFEM_FOREACH_THREAD(qx,x,Q1D)
{
MFEM_FOREACH_THREAD(qy,y,Q1D)
{
const real_t detJ = J(qx,qy,e);
const real_t coeff = const_c ? C(0,0,0) : C(qx,qy,e);
v(qx,qy,e) = W(qx,qy) * coeff * (by_val ? detJ : 1.0/detJ);
}
}
});
}
else
{
MFEM_ABORT("Not supported.");
}
}
void MassIntegrator::AssembleDiagonalPA(Vector &diag)
@@ -139,6 +172,8 @@ void MassIntegrator::AssembleDiagonalPA(Vector &diag)
void MassIntegrator::AddMultPA(const Vector &x, Vector &y) const
{
MFEM_PERF_FUNCTION;
if (DeviceCanUseCeed())
{
ceedOp->AddMult(x, y);
+29 -29
View File
@@ -346,13 +346,13 @@ private:
template<typename T>
T operator() (const blitz::TinyVector<T,3>& x) const
{
const int el_order = el->GetOrder();
std::vector<T> u1(el_order+1);
std::vector<T> u2(el_order+1);
std::vector<T> u3(el_order+1);
TmplPoly_1D::CalcBernstein(el_order, x[0], u1.data());
TmplPoly_1D::CalcBernstein(el_order, x[1], u2.data());
TmplPoly_1D::CalcBernstein(el_order, x[2], u3.data());
int el_order=el->GetOrder();
T u1[el_order+1];
T u2[el_order+1];
T u3[el_order+1];
TmplPoly_1D::CalcBernstein(el_order, x[0], u1);
TmplPoly_1D::CalcBernstein(el_order, x[1], u2);
TmplPoly_1D::CalcBernstein(el_order, x[2], u3);
const Array<int>& dof_map=el->GetDofMap();
@@ -370,17 +370,17 @@ private:
template<typename T>
blitz::TinyVector<T,3> grad(const blitz::TinyVector<T,3>& x) const
{
const int el_order = el->GetOrder();
std::vector<T> u1(el_order+1);
std::vector<T> u2(el_order+1);
std::vector<T> u3(el_order+1);
std::vector<T> d1(el_order+1);
std::vector<T> d2(el_order+1);
std::vector<T> d3(el_order+1);
int el_order=el->GetOrder();
T u1[el_order+1];
T u2[el_order+1];
T u3[el_order+1];
T d1[el_order+1];
T d2[el_order+1];
T d3[el_order+1];
TmplPoly_1D::CalcBernstein(el_order,x[0], u1.data(), d1.data());
TmplPoly_1D::CalcBernstein(el_order,x[1], u2.data(), d2.data());
TmplPoly_1D::CalcBernstein(el_order,x[2], u3.data(), d3.data());
TmplPoly_1D::CalcBernstein(el_order,x[0], u1, d1);
TmplPoly_1D::CalcBernstein(el_order,x[1], u2, d2);
TmplPoly_1D::CalcBernstein(el_order,x[2], u3, d3);
blitz::TinyVector<T,3> res(T(0.0),T(0.0),T(0.0));
@@ -415,11 +415,11 @@ private:
template<typename T>
T operator() (const blitz::TinyVector<T,2>& x) const
{
const int el_order = el->GetOrder();
std::vector<T> u1(el_order+1);
std::vector<T> u2(el_order+1);
TmplPoly_1D::CalcBernstein(el_order, x[0], u1.data());
TmplPoly_1D::CalcBernstein(el_order, x[1], u2.data());
int el_order=el->GetOrder();
T u1[el_order+1];
T u2[el_order+1];
TmplPoly_1D::CalcBernstein(el_order, x[0], u1);
TmplPoly_1D::CalcBernstein(el_order, x[1], u2);
const Array<int>& dof_map=el->GetDofMap();
@@ -437,14 +437,14 @@ private:
template<typename T>
blitz::TinyVector<T,2> grad(const blitz::TinyVector<T,2>& x) const
{
const int el_order = el->GetOrder();
std::vector<T> u1(el_order+1);
std::vector<T> u2(el_order+1);
std::vector<T> d1(el_order+1);
std::vector<T> d2(el_order+1);
int el_order=el->GetOrder();
T u1[el_order+1];
T u2[el_order+1];
T d1[el_order+1];
T d2[el_order+1];
TmplPoly_1D::CalcBernstein(el_order,x[0], u1.data(), d1.data());
TmplPoly_1D::CalcBernstein(el_order,x[1], u2.data(), d2.data());
TmplPoly_1D::CalcBernstein(el_order,x[0], u1, d1);
TmplPoly_1D::CalcBernstein(el_order,x[1], u2, d2);
blitz::TinyVector<T,2> res(T(0.0),T(0.0));
+32 -10
View File
@@ -51,7 +51,7 @@ void LinearFormExtension::Assemble()
{
// scan the attributes to set the markers to 0 or 1
const int NE = fes.GetNE();
const auto attr = attributes->Read();
const auto attr = attributes.Read();
const auto dimk = domain_integs_marker_k->Read();
auto markers_w = markers.Write();
mfem::forall(NE, [=] MFEM_HOST_DEVICE (int e)
@@ -93,14 +93,13 @@ void LinearFormExtension::Assemble()
else
{
// scan the attributes to set the markers to 0 or 1
const int NBE = bdr_face_attributes->Size();
const auto attr = bdr_face_attributes->Read();
const int NBE = bdr_attributes.Size();
const auto attr = bdr_attributes.Read();
const auto attr_markers = boundary_integs_marker_k->Read();
auto markers_w = bdr_markers.Write();
mfem::forall(NBE, [=] MFEM_HOST_DEVICE(int e)
mfem::forall(NBE, [=] MFEM_HOST_DEVICE (int e)
{
markers_w[e] =
attr[e] > 0 ? (attr_markers[attr[e] - 1] == 1) : false;
markers_w[e] = attr_markers[attr[e]-1] == 1;
});
}
@@ -126,7 +125,8 @@ void LinearFormExtension::Update()
//markers.UseDevice(true);
// Gather the attributes on the host from all the elements
attributes = &mesh.GetElementAttributes();
attributes.SetSize(NE);
for (int i = 0; i < NE; ++i) { attributes[i] = mesh.GetAttribute(i); }
elem_restrict_lex = fes.GetElementRestriction(ordering);
MFEM_VERIFY(elem_restrict_lex, "Element restriction not available");
@@ -136,12 +136,34 @@ void LinearFormExtension::Update()
if (lf->boundary_integs.Size() > 0)
{
bdr_face_attributes = &mesh.GetBdrFaceAttributes();
const int nf_bdr = bdr_face_attributes->Size();
const int nf_bdr = fes.GetNFbyType(FaceType::Boundary);
bdr_markers.SetSize(nf_bdr);
// bdr_markers.UseDevice(true);
// The face restriction will give us "face E-vectors" on the boundary that
// are numbered in the order of the faces of mesh. This numbering will be
// different than the numbering of the boundary elements. We compute
// mappings so that the array `bdr_attributes[i]` gives the boundary
// attribute of the `i`th boundary face in the mesh face order.
std::unordered_map<int,int> f_to_be;
for (int i = 0; i < mesh.GetNBE(); ++i)
{
const int f = mesh.GetBdrElementFaceIndex(i);
f_to_be[f] = i;
}
MFEM_VERIFY(size_t(nf_bdr) == f_to_be.size(), "Incompatible sizes");
bdr_attributes.SetSize(nf_bdr);
int f_ind = 0;
for (int f = 0; f < mesh.GetNumFaces(); ++f)
{
if (f_to_be.find(f) != f_to_be.end())
{
const int be = f_to_be[f];
bdr_attributes[f_ind] = mesh.GetBdrAttribute(be);
++f_ind;
}
}
bdr_restrict_lex =
dynamic_cast<const FaceRestriction*>(
fes.GetFaceRestriction(ordering, FaceType::Boundary,
+2 -2
View File
@@ -14,6 +14,7 @@
#include "../general/array.hpp"
#include "../linalg/vector.hpp"
#include "fespace.hpp"
namespace mfem
{
@@ -25,8 +26,7 @@ class LinearForm;
class LinearFormExtension
{
/// Attributes of all mesh elements.
const Array<int> *attributes; // Not owned
const Array<int> *bdr_face_attributes; // Not owned
Array<int> attributes, bdr_attributes;
/// Temporary markers for device kernels.
Array<int> markers, bdr_markers;
+2
View File
@@ -365,6 +365,8 @@ FiniteElementSpace &LORBase::GetFESpace() const
void LORBase::AssembleSystem(BilinearForm &a_ho, const Array<int> &ess_dofs)
{
MFEM_PERF_FUNCTION;
A.Clear();
delete a;
if (BatchedLORAssembly::FormIsSupported(a_ho))
+4
View File
@@ -360,6 +360,8 @@ void BatchedLORAssembly::FillJAndData(SparseMatrix &A) const
void BatchedLORAssembly::SparseIJToCSR(OperatorHandle &A) const
{
MFEM_PERF_FUNCTION;
const int nvdof = fes_ho.GetVSize();
// If A contains an existing SparseMatrix, reuse it (and try to reuse its
@@ -417,6 +419,8 @@ static void Assemble_(LOR_KERNEL &kernel, int dim, int sdim, int order)
template <typename LOR_KERNEL>
void BatchedLORAssembly::AssemblyKernel(BilinearForm &a)
{
MFEM_PERF_FUNCTION;
LOR_KERNEL kernel(a, fes_ho, X_vert, sparse_ij, sparse_mapping);
const int dim = fes_ho.GetMesh()->Dimension();
+2
View File
@@ -184,6 +184,8 @@ void BatchedLOR_H1::Assemble2D()
template <int ORDER>
void BatchedLOR_H1::Assemble3D()
{
MFEM_PERF_FUNCTION;
const int nel_ho = fes_ho.GetNE();
static constexpr int nv = 8;
static constexpr int dim = 3;
+136 -58
View File
@@ -17,7 +17,10 @@ namespace mfem
MultigridBase::MultigridBase()
: cycleType(CycleType::VCYCLE), preSmoothingSteps(1), postSmoothingSteps(1),
nrhs(0)
{}
{
coarse_solver = nullptr;
own_coarse_solver = false;
}
MultigridBase::MultigridBase(const Array<Operator*>& operators_,
const Array<Solver*>& smoothers_,
@@ -29,12 +32,18 @@ MultigridBase::MultigridBase(const Array<Operator*>& operators_,
{
operators_.Copy(operators);
smoothers_.Copy(smoothers);
coarse_solver = nullptr;
ownedOperators_.Copy(ownedOperators);
ownedSmoothers_.Copy(ownedSmoothers);
own_coarse_solver = false;
}
MultigridBase::~MultigridBase()
{
if (own_coarse_solver)
{
delete coarse_solver;
}
for (int i = 0; i < operators.Size(); ++i)
{
if (ownedOperators[i])
@@ -56,16 +65,17 @@ void MultigridBase::InitVectors() const
X.SetSize(M, nrhs);
Y.SetSize(M, nrhs);
R.SetSize(M, nrhs);
Z.SetSize(M, nrhs);
for (int i = 0; i < X.NumRows(); ++i)
for (int i = 0; i < M; ++i)
{
const int n = operators[i]->Height();
for (int j = 0; j < X.NumCols(); ++j)
for (int j = 0; j < nrhs; ++j)
{
X(i, j) = new Vector(n);
Y(i, j) = new Vector(n);
if (i < M - 1)
{
X(i, j) = new Vector(n);
Y(i, j) = new Vector(n);
}
R(i, j) = new Vector(n);
Z(i, j) = new Vector(n);
}
}
}
@@ -76,10 +86,12 @@ void MultigridBase::EraseVectors() const
{
for (int j = 0; j < X.NumCols(); ++j)
{
delete X(i, j);
delete Y(i, j);
if (i < X.NumRows() - 1)
{
delete X(i, j);
delete Y(i, j);
}
delete R(i, j);
delete Z(i, j);
}
}
}
@@ -95,6 +107,12 @@ void MultigridBase::AddLevel(Operator* op, Solver* smoother,
ownedSmoothers.Append(ownSmoother);
}
void MultigridBase::AddCoarseSolver(Solver *c_solver, bool own_c_solver)
{
coarse_solver = c_solver;
own_coarse_solver = own_c_solver;
}
void MultigridBase::SetCycleType(CycleType cycleType_, int preSmoothingSteps_,
int postSmoothingSteps_)
{
@@ -105,10 +123,12 @@ void MultigridBase::SetCycleType(CycleType cycleType_, int preSmoothingSteps_,
void MultigridBase::Mult(const Vector& x, Vector& y) const
{
Array<const Vector*> X_(1);
Array<Vector*> Y_(1);
X_[0] = &x;
Y_[0] = &y;
const Vector *x_array[1] = { &x };
Array<const Vector*> X_(x_array, 1); // no heap allocation
Vector *y_array[1] = { &y };
Array<Vector*> Y_(y_array, 1); // no heap allocation
ArrayMult(X_, Y_);
}
@@ -119,11 +139,6 @@ void MultigridBase::ArrayMult(const Array<const Vector*>& X_,
"Multigrid solver does not have operators set!");
MFEM_ASSERT(X_.Size() == Y_.Size(),
"Number of columns mismatch in MultigridBase::Mult!");
if (iterative_mode)
{
MFEM_WARNING("Multigrid solver does not use iterative_mode and ignores "
"the initial guess!");
}
// Add capacity as necessary
nrhs = X_.Size();
@@ -134,96 +149,159 @@ void MultigridBase::ArrayMult(const Array<const Vector*>& X_,
for (int j = 0; j < nrhs; ++j)
{
MFEM_ASSERT(X_[j] && Y_[j], "Missing Vector in MultigridBase::Mult!");
*X(M - 1, j) = *X_[j];
*Y(M - 1, j) = 0.0;
}
Cycle(M - 1);
for (int j = 0; j < nrhs; ++j)
{
*Y_[j] = *Y(M - 1, j);
X(M - 1, j) = const_cast<Vector*>(X_[j]);
Y(M - 1, j) = Y_[j];
}
const bool zero = !iterative_mode;
Cycle(M - 1, zero);
}
void MultigridBase::SmoothingStep(int level, bool zero, bool transpose) const
{
// y = y + S (x - A y) or y = y + S^T (x - A y)
// Note: 'zero' == true means that Y(level,*) are not initialized and we
// should assume that the input they typically provide to this call is zeros.
// We can't use the smoothers' iterative mode since we don't know if they
// actually support it, so we always turn the iterative mode off to properly
// use smoothers that do support it.
smoothers[level]->iterative_mode = false;
if (zero)
{
Array<Vector *> X_(X[level], nrhs), Y_(Y[level], nrhs);
GetSmootherAtLevel(level)->ArrayMult(X_, Y_);
MFEM_ASSERT(!transpose, "internal error!");
const Array<const Vector *> cX_((const Vector **)(X[level]), nrhs);
Array<Vector *> Y_(Y[level], nrhs);
GetSmootherAtLevel(level)->ArrayMult(cX_, Y_);
}
else
{
Array<Vector *> Y_(Y[level], nrhs), R_(R[level], nrhs),
Z_(Z[level], nrhs);
const Array<const Vector *> cY_((const Vector **)(Y[level]), nrhs),
cR_((const Vector **)(R[level]), nrhs);
Array<Vector *> Y_(Y[level], nrhs), R_(R[level], nrhs);
GetOperatorAtLevel(level)->ArrayMult(cY_, R_);
for (int j = 0; j < nrhs; ++j)
{
*R_[j] = *X(level, j);
// *R_[j] = *X(level, j) - *R_[j]
subtract(*X(level, j), *R_[j], *R_[j]);
}
GetOperatorAtLevel(level)->ArrayAddMult(Y_, R_, -1.0);
if (transpose)
{
GetSmootherAtLevel(level)->ArrayMultTranspose(R_, Z_);
GetSmootherAtLevel(level)->ArrayAddMultTranspose(cR_, Y_);
}
else
{
GetSmootherAtLevel(level)->ArrayMult(R_, Z_);
}
for (int j = 0; j < nrhs; ++j)
{
*Y_[j] += *Z_[j];
GetSmootherAtLevel(level)->ArrayAddMult(cR_, Y_);
}
}
}
void MultigridBase::Cycle(int level) const
void MultigridBase::CoarseSolve(bool zero) const
{
// Coarse solve
if (level == 0)
// See the comment about iterative mode in SmoothingStep()
coarse_solver->iterative_mode = false;
if (zero)
{
SmoothingStep(0, true, false);
const Array<const Vector *> cX_((const Vector **)(X[0]), nrhs);
Array<Vector *> Y_(Y[0], nrhs);
coarse_solver->ArrayMult(cX_, Y_);
}
else
{
const Array<const Vector *> cY_((const Vector **)(Y[0]), nrhs),
cR_((const Vector **)(R[0]), nrhs);
Array<Vector *> Y_(Y[0], nrhs), R_(R[0], nrhs);
GetOperatorAtLevel(0)->ArrayMult(cY_, R_);
for (int j = 0; j < nrhs; ++j)
{
// *R_[j] = *X(0, j) - *R_[j]
subtract(*X(0, j), *R_[j], *R_[j]);
}
coarse_solver->ArrayAddMult(cR_, Y_);
}
}
void MultigridBase::Cycle(int level, bool zero) const
{
// Note: 'zero' == true means that Y(level,*) are not initialized and we
// should assume that the input they typically provide to this call is zeros.
// Coarse solve
if (level == 0 && !coarse_solver)
{
SmoothingStep(0, zero, false);
return;
}
// Pre-smooth
for (int i = 0; i < preSmoothingSteps; ++i)
{
SmoothingStep(level, (cycleType == CycleType::VCYCLE && i == 0), false);
SmoothingStep(level, zero && (i == 0), false);
}
// Coarse solve with 'coarse_solver'
if (level == 0)
{
CoarseSolve(preSmoothingSteps == 0 && zero);
goto mg_post_smooth;
}
// Compute residual and restrict
if (preSmoothingSteps == 0 && zero)
{
Array<Vector *> Y_(Y[level], nrhs), R_(R[level], nrhs),
X_(X[level - 1], nrhs);
const Array<const Vector *> cX_l((const Vector **)(X[level]), nrhs);
Array<Vector *> X_lm1(X[level - 1], nrhs);
GetProlongationAtLevel(level - 1)->ArrayMultTranspose(cX_l, X_lm1);
}
else
{
const Array<const Vector *> cY_((const Vector **)(Y[level]), nrhs),
cR_((const Vector **)(R[level]), nrhs);
Array<Vector *> R_(R[level], nrhs), X_(X[level - 1], nrhs);
GetOperatorAtLevel(level)->ArrayMult(cY_, R_);
for (int j = 0; j < nrhs; ++j)
{
*R_[j] = *X(level, j);
}
GetOperatorAtLevel(level)->ArrayAddMult(Y_, R_, -1.0);
GetProlongationAtLevel(level - 1)->ArrayMultTranspose(R_, X_);
for (int j = 0; j < nrhs; ++j)
{
*Y(level - 1, j) = 0.0;
// *R_[j] = *X(level, j) - *R_[j]
subtract(*X(level, j), *R_[j], *R_[j]);
}
GetProlongationAtLevel(level - 1)->ArrayMultTranspose(cR_, X_);
}
// Corrections
Cycle(level - 1);
Cycle(level - 1, true);
if (cycleType == CycleType::WCYCLE)
{
Cycle(level - 1);
// If the coarse solve at level 0 is "exact" solve, then we don't want to
// repeat it.
// To support multiple level 0 coarse-grid corrections, one can wrap that
// smoother in an SLI solver and use that instead.
if (level > 1) { Cycle(level - 1, false); }
}
// Prolongate and add
{
Array<Vector *> Y_(Y[level - 1], nrhs), Z_(Z[level], nrhs);
GetProlongationAtLevel(level - 1)->ArrayMult(Y_, Z_);
for (int j = 0; j < nrhs; ++j)
const Array<const Vector *> cY_lm1((const Vector **)(Y[level - 1]), nrhs);
Array<Vector *> Y_l(Y[level], nrhs);
if (preSmoothingSteps == 0 && zero)
{
*Y(level, j) += *Z_[j];
GetProlongationAtLevel(level - 1)->ArrayMult(cY_lm1, Y_l);
}
else
{
GetProlongationAtLevel(level - 1)->ArrayAddMult(cY_lm1, Y_l);
}
}
mg_post_smooth:
// Post-smooth
for (int i = 0; i < postSmoothingSteps; ++i)
{
+20 -2
View File
@@ -36,12 +36,14 @@ protected:
Array<Solver*> smoothers;
Array<bool> ownedOperators;
Array<bool> ownedSmoothers;
Solver *coarse_solver; /// can be NULL, see AddCoarseSolver()
bool own_coarse_solver;
CycleType cycleType;
int preSmoothingSteps;
int postSmoothingSteps;
mutable Array2D<Vector*> X, Y, R, Z;
mutable Array2D<Vector*> X, Y, R;
mutable int nrhs;
public:
@@ -65,6 +67,16 @@ public:
void AddLevel(Operator* op, Solver* smoother, bool ownOperator,
bool ownSmoother);
/// Adds a coarse solver for level 0 to work in tandem with the smoother
/** If this coarse solver is not given, the smoother at level 0 is used as
the coarse solver. When this coarse solver is given, the smoother at
level 0 is used similar to the smoothers at other levels. Thus, the
action at level 0 consists of:
- pre-smoothing steps with smoother 0,
- solve step with @a c_solver,
- post-smoothing steps with smoother 0. */
void AddCoarseSolver(Solver *c_solver, bool own_c_solver);
/// Returns the number of levels
int NumLevels() const { return operators.Size(); }
@@ -118,11 +130,14 @@ public:
private:
/// Application of a multigrid cycle at particular level
void Cycle(int level) const;
void Cycle(int level, bool zero) const;
/// Application of a pre-/post-smoothing step at particular level
void SmoothingStep(int level, bool zero, bool transpose) const;
/// Perform a coarse solve with 'coarse_solve' (must be non-NULL)
void CoarseSolve(bool zero) const;
/// Allocate or destroy temporary storage
void InitVectors() const;
void EraseVectors() const;
@@ -202,6 +217,9 @@ public:
/// Recover the solution of a linear system formed with FormFineLinearSystem()
void RecoverFineFEMSolution(const Vector& X, const Vector& b, Vector& x);
const Array<int> &GetFineEssentialTrueDofs() const
{ return *essentialTrueDofs.Last(); }
};
} // namespace mfem
+2
View File
@@ -124,6 +124,8 @@ void ParBilinearForm::pAllocMat()
void ParBilinearForm::ParallelRAP(SparseMatrix &loc_A, OperatorHandle &A,
bool steal_loc_A)
{
MFEM_PERF_FUNCTION;
ParFiniteElementSpace &pfespace = *ParFESpace();
// Create a block diagonal parallel matrix
-591
View File
@@ -1,591 +0,0 @@
// Copyright (c) 2010-2025, Lawrence Livermore National Security, LLC. Produced
// at the Lawrence Livermore National Laboratory. All Rights reserved. See files
// LICENSE and NOTICE for details. LLNL-CODE-806117.
//
// This file is part of the MFEM library. For more information and source code
// availability visit https://mfem.org.
//
// MFEM is free software; you can redistribute it and/or modify it under the
// terms of the BSD-3 license. We welcome feedback and contributions, see file
// CONTRIBUTING.md for details.
#include "pderefmat_op.hpp"
#ifdef MFEM_USE_MPI
#include "fes_kernels.hpp"
/// \cond DO_NOT_DOCUMENT
namespace mfem
{
namespace internal
{
template <Ordering::Type Order, bool Atomic>
static void ParDerefMultKernelImpl(const ParDerefineMatrixOp &op,
const Vector &x, Vector &y)
{
// pack sends
if (op.xghost_send.Size())
{
auto src = x.Read();
auto idcs = op.send_permutations.Read();
auto dst = Device::GetGPUAwareMPI() ? op.xghost_send.Write()
: op.xghost_send.HostWrite();
auto vdims = op.fespace->GetVDim();
auto sptr = op.send_segment_idcs.Read();
auto lptr = op.send_segments.Read();
auto old_ndofs = x.Size() / vdims;
forall(op.send_permutations.Size(), [=] MFEM_HOST_DEVICE(int i)
{
int seg = sptr[i];
int width = lptr[seg + 1] - lptr[seg];
auto tdst = dst + i + lptr[seg] * vdims;
int sign = 1;
int col = idcs[i];
if (col < 0)
{
sign = -1;
col = -1 - col;
}
for (int vdim = 0; vdim < vdims; ++vdim)
{
tdst[vdim * width] =
sign
* src[Order == Ordering::byNODES ? (col + vdim * old_ndofs)
: (col * vdims + vdim)];
}
});
// TODO: is this needed so we can send the packed data correctly?
// unclear for GPU-aware MPI, definitely required otherwise
MFEM_DEVICE_SYNC;
}
// initialize off-diagonal receive and send
op.requests.clear();
if (op.xghost_recv.Size())
{
auto vdims = op.fespace->GetVDim();
auto rcv = Device::GetGPUAwareMPI() ? op.xghost_recv.Write()
: op.xghost_recv.HostWrite();
for (int i = 0; i < op.recv_ranks.Size(); ++i)
{
op.requests.emplace_back();
MPI_Irecv(rcv + op.recv_segments[i] * vdims,
(op.recv_segments[i + 1] - op.recv_segments[i]) * vdims,
MPITypeMap<real_t>::mpi_type, op.recv_ranks[i],
MessageTag::DEREFINEMENT_MATRIX_CONSTRUCTION_DATA,
op.fespace->GetComm(), &op.requests.back());
}
}
if (op.xghost_send.Size())
{
auto vdims = op.fespace->GetVDim();
// only is a GPU mem ptr if GPU-aware MPI is enabled
auto dst = Device::GetGPUAwareMPI() ? op.xghost_send.Write()
: op.xghost_send.HostWrite();
for (int i = 0; i < op.send_ranks.Size(); ++i)
{
op.requests.emplace_back();
MPI_Isend(dst + op.send_segments[i] * vdims,
(op.send_segments[i + 1] - op.send_segments[i]) * vdims,
MPITypeMap<real_t>::mpi_type, op.send_ranks[i],
MessageTag::DEREFINEMENT_MATRIX_CONSTRUCTION_DATA,
op.fespace->GetComm(), &op.requests.back());
}
}
{
// diagonal
DerefineMatrixOpMultFunctor<Order, Atomic, true> func;
func.xptr = x.Read();
y.UseDevice();
y = 0.;
func.yptr = y.ReadWrite();
func.bsptr = op.block_storage.Read();
func.boptr = op.block_offsets.Read();
func.brptr = op.block_row_idcs_offsets.Read();
func.bcptr = op.block_col_idcs_offsets.Read();
func.rptr = op.row_idcs.Read();
func.cptr = op.col_idcs.Read();
func.vdims = op.fespace->GetVDim();
func.nblocks = op.block_offsets.Size();
func.width = op.Width() / func.vdims;
func.height = op.Height() / func.vdims;
func.Run(op.max_rows);
}
// wait for comm to finish, if any
if (op.requests.size())
{
MPI_Waitall(op.requests.size(), op.requests.data(), MPI_STATUSES_IGNORE);
if (op.xghost_recv.Size())
{
// off-diagonal kernel
DerefineMatrixOpMultFunctor<Order, Atomic, false> func;
// directly read from host-pinned memory if not using GPU-aware MPI
func.xptr = Device::GetGPUAwareMPI() ? op.xghost_recv.Read()
: op.xghost_recv.HostRead();
func.yptr = y.ReadWrite();
func.bsptr = op.block_storage.Read();
func.boptr = op.off_diag_block_offsets.Read();
func.brptr = op.block_off_diag_row_idcs_offsets.Read();
func.rsptr = op.recv_segment_idcs.Read();
func.segptr = op.recv_segments.Read();
func.coptr = op.block_off_diag_col_offsets.Read();
func.bwptr = op.block_off_diag_widths.Read();
func.rptr = op.row_off_diag_idcs.Read();
func.vdims = op.fespace->GetVDim();
func.nblocks = op.off_diag_block_offsets.Size();
func.width = op.xghost_recv.Size() / func.vdims;
func.height = op.Height() / func.vdims;
func.Run(op.max_rows);
}
}
}
} // namespace internal
template <Ordering::Type Order, bool Atomic>
ParDerefineMatrixOp::MultKernelType ParDerefineMatrixOp::MultKernel::Kernel()
{
return internal::ParDerefMultKernelImpl<Order, Atomic>;
}
ParDerefineMatrixOp::MultKernelType
ParDerefineMatrixOp::MultKernel::Fallback(Ordering::Type, bool)
{
MFEM_ABORT("invalid MultKernel parameters");
}
ParDerefineMatrixOp::Kernels::Kernels()
{
MultKernel::Specialization<Ordering::byNODES, false>::Add();
MultKernel::Specialization<Ordering::byVDIM, false>::Add();
MultKernel::Specialization<Ordering::byNODES, true>::Add();
MultKernel::Specialization<Ordering::byVDIM, true>::Add();
}
void ParDerefineMatrixOp::Mult(const Vector &x, Vector &y) const
{
const bool is_dg = fespace->FEColl()->GetContType()
== FiniteElementCollection::DISCONTINUOUS;
// DG needs atomic summation
MultKernel::Run(fespace->GetOrdering(), is_dg, *this, x, y);
// use this to prevent xghost* from being re-purposed for subsequent Mult
// calls
MFEM_DEVICE_SYNC;
}
ParDerefineMatrixOp::ParDerefineMatrixOp(ParFiniteElementSpace &fespace_,
int old_ndofs,
const Table *old_elem_dof,
const Table *old_elem_fos)
: Operator(fespace_.GetVSize(), old_ndofs * fespace_.GetVDim()),
fespace(&fespace_)
{
static Kernels kernels;
constexpr int max_team_size = 256;
const int NRanks = fespace->GetNRanks();
const int nrk = HYPRE_AssumedPartitionCheck() ? 2 : NRanks;
MFEM_VERIFY(fespace->Nonconforming(),
"Not implemented for conforming meshes.");
MFEM_VERIFY(fespace->old_dof_offsets[nrk],
"Missing previous (finer) space.");
const int MyRank = fespace->GetMyRank();
ParNCMesh *old_pncmesh = fespace->GetParMesh()->pncmesh;
const CoarseFineTransformations &dtrans =
old_pncmesh->GetDerefinementTransforms();
const Array<int> &old_ranks = old_pncmesh->GetDerefineOldRanks();
const bool is_dg = fespace->FEColl()->GetContType()
== FiniteElementCollection::DISCONTINUOUS;
DenseMatrix localRVO; // for variable-order only
DenseTensor localR[Geometry::NumGeom];
int diag_rows = 0;
int off_diag_rows = 0;
int diag_cols = 0;
auto get_ldofs = [&](int k) -> int
{
const Embedding &emb = dtrans.embeddings[k];
if (fespace->IsVariableOrder())
{
const FiniteElement *fe = fespace->GetFE(emb.parent);
return fe->GetDof();
}
else
{
Geometry::Type geom =
fespace->GetParMesh()->GetElementBaseGeometry(emb.parent);
return fespace->FEColl()->FiniteElementForGeometry(geom)->GetDof();
}
};
Array<int> dofs, old_dofs;
max_rows = 1;
// first pass:
// - determine memory block lengths
// - identify dofs in x we need to send/receive
// don't need to send the indices, fine rank will re-arrange and sign
// change x before transmitting the ghost data
// key: coarse rank to send to
// value: old dofs to send (with sign)
std::map<int, std::vector<int>> to_send;
// key: fine rank
// value: indices into dtrans.embeddings
std::map<int, std::vector<int>> od_ks;
// key: fine rank
// value: recv segment length
std::map<int, int> od_seg_lens;
int send_len = 0;
int recv_len = 0;
// size of block_storage, if fespace->IsVariableOrder()
// otherwise unused
int total_size = 0;
int num_diagonal_blocks = 0;
int num_offdiagonal_blocks = 0;
for (int k = 0; k < dtrans.embeddings.Size(); ++k)
{
const Embedding &emb = dtrans.embeddings[k];
int fine_rank = old_ranks[k];
int coarse_rank = (emb.parent < 0) ? (-1 - emb.parent)
: old_pncmesh->ElementRank(emb.parent);
if (coarse_rank != MyRank && fine_rank == MyRank)
{
// this rank needs to send data in x to course_rank
old_elem_dof->GetRow(k, old_dofs);
auto &tmp = to_send[coarse_rank];
send_len += old_dofs.Size();
for (int i = 0; i < old_dofs.Size(); ++i)
{
tmp.emplace_back(old_dofs[i]);
}
}
else if (coarse_rank == MyRank && fine_rank != MyRank)
{
// this rank needs to receive data in x from fine_rank
MFEM_ASSERT(emb.parent >= 0, "");
auto ldofs = get_ldofs(k);
off_diag_rows += ldofs;
recv_len += ldofs;
od_ks[fine_rank].emplace_back(k);
od_seg_lens[fine_rank] += ldofs;
++num_offdiagonal_blocks;
if (fespace->IsVariableOrder())
{
total_size += ldofs * ldofs;
}
}
else if (coarse_rank == MyRank && fine_rank == MyRank)
{
MFEM_ASSERT(emb.parent >= 0, "");
// diagonal
++num_diagonal_blocks;
auto ldofs = get_ldofs(k);
diag_rows += ldofs;
diag_cols += ldofs;
if (fespace->IsVariableOrder())
{
total_size += ldofs * ldofs;
}
}
}
send_segments.SetSize(to_send.size() + 1);
send_segments.HostWrite();
send_ranks.SetSize(to_send.size());
send_ranks.HostWrite();
{
int idx = 0;
send_segments[0] = 0;
for (auto &tmp : to_send)
{
send_ranks[idx] = tmp.first;
send_segments[idx + 1] = send_segments[idx] + tmp.second.size();
++idx;
}
}
recv_segment_idcs.SetSize(off_diag_rows);
recv_segment_idcs.HostWrite();
recv_segments.SetSize(od_ks.size() + 1);
recv_segments.HostWrite();
recv_ranks.SetSize(od_ks.size());
recv_ranks.HostWrite();
// set sizes
row_idcs.SetSize(diag_rows);
row_idcs.HostWrite();
row_off_diag_idcs.SetSize(off_diag_rows);
row_off_diag_idcs.HostWrite();
col_idcs.SetSize(diag_cols);
col_idcs.HostWrite();
block_row_idcs_offsets.SetSize(num_diagonal_blocks + 1);
block_row_idcs_offsets.HostWrite();
block_col_idcs_offsets.SetSize(num_diagonal_blocks + 1);
block_col_idcs_offsets.HostWrite();
block_off_diag_row_idcs_offsets.SetSize(num_offdiagonal_blocks + 1);
block_off_diag_row_idcs_offsets.HostWrite();
block_off_diag_col_offsets.SetSize(num_offdiagonal_blocks);
block_off_diag_col_offsets.HostWrite();
block_off_diag_widths.SetSize(num_offdiagonal_blocks);
block_off_diag_widths.HostWrite();
pack_col_idcs.SetSize(send_len);
// memory manager doesn't appear to have a graceful fallback for
// HOST_PINNED if not built with CUDA or HIP
#if defined(MFEM_USE_CUDA) or defined(MFEM_USE_HIP)
xghost_send.SetSize(send_len * fespace->GetVDim(),
Device::GetGPUAwareMPI() ? MemoryType::DEFAULT
: MemoryType::HOST_PINNED);
xghost_recv.SetSize(recv_len * fespace->GetVDim(),
Device::GetGPUAwareMPI() ? MemoryType::DEFAULT
: MemoryType::HOST_PINNED);
#else
xghost_send.SetSize(send_len * fespace->GetVDim());
xghost_recv.SetSize(recv_len * fespace->GetVDim());
#endif
send_permutations.SetSize(send_len);
send_segment_idcs.SetSize(send_len);
block_offsets.SetSize(num_diagonal_blocks);
block_offsets.HostWrite();
off_diag_block_offsets.SetSize(num_offdiagonal_blocks);
off_diag_block_offsets.HostWrite();
int geom_offsets[Geometry::NumGeom];
real_t *bs_ptr;
if (fespace->IsVariableOrder())
{
block_storage.SetSize(total_size);
bs_ptr = block_storage.HostWrite();
// compute block data later
}
else
{
// compression scheme:
// block_offsets is the start of each block, potentially repeated
// only need to store localR for used shapes
Mesh::GeometryList elem_geoms(*fespace->GetMesh());
int size = 0;
for (int i = 0; i < elem_geoms.Size(); ++i)
{
fespace->GetLocalDerefinementMatrices(elem_geoms[i],
localR[elem_geoms[i]]);
geom_offsets[elem_geoms[i]] = size;
size += localR[elem_geoms[i]].TotalSize();
}
block_storage.SetSize(size);
bs_ptr = block_storage.HostWrite();
// copy blocks into block_storage
for (int i = 0; i < elem_geoms.Size(); ++i)
{
std::copy(localR[elem_geoms[i]].Data(),
localR[elem_geoms[i]].Data()
+ localR[elem_geoms[i]].TotalSize(),
bs_ptr);
bs_ptr += localR[elem_geoms[i]].TotalSize();
}
}
// second pass:
// - initialize buffers
{
auto ptr = send_permutations.HostWrite();
auto ptr2 = send_segment_idcs.HostWrite();
int i = 0;
for (auto &v : to_send)
{
ptr = std::copy(v.second.begin(), v.second.end(), ptr);
for (size_t idx = 0; idx < v.second.size(); ++idx)
{
*ptr2 = i;
++ptr2;
}
++i;
}
}
block_row_idcs_offsets[0] = 0;
block_col_idcs_offsets[0] = 0;
block_off_diag_row_idcs_offsets[0] = 0;
Array<int> mark(fespace->GetNDofs());
mark = 0;
{
int idx = 0;
recv_segments[0] = 0;
for (auto &v : od_seg_lens)
{
recv_ranks[idx] = v.first;
recv_segments[idx + 1] = recv_segments[idx] + v.second;
++idx;
}
}
// key: index into dtrans.embeddings
// value: off-diagonal block offset, od_ridx, seg id
std::unordered_map<int, std::array<int, 3>> ks_map;
{
int od_ridx = 0;
int seg_id = 0;
for (auto &v1 : od_ks)
{
for (auto k : v1.second)
{
auto &tmp = ks_map[k];
tmp[0] = ks_map.size() - 1;
tmp[1] = od_ridx;
tmp[2] = seg_id;
od_ridx += get_ldofs(k);
}
++seg_id;
}
}
int diag_idx = 0;
int var_offset = 0;
int ridx = 0;
int cidx = 0;
// can't break this up into separate diagonals/off-diagonals loops because
// of mark
for (int k = 0; k < dtrans.embeddings.Size(); ++k)
{
const Embedding &emb = dtrans.embeddings[k];
if (emb.parent < 0)
{
continue;
}
int fine_rank = old_ranks[k];
int coarse_rank = (emb.parent < 0) ? (-1 - emb.parent)
: old_pncmesh->ElementRank(emb.parent);
if (coarse_rank == MyRank)
{
// either diagonal or off-diagonal
Geometry::Type geom =
fespace->GetMesh()->GetElementBaseGeometry(emb.parent);
if (fespace->IsVariableOrder())
{
const FiniteElement *fe = fespace->GetFE(emb.parent);
const DenseTensor &pmats = dtrans.point_matrices[geom];
const int ldof = fe->GetDof();
IsoparametricTransformation isotr;
isotr.SetIdentityTransformation(geom);
localRVO.SetSize(ldof, ldof);
isotr.SetPointMat(pmats(emb.matrix));
// Local restriction is size ldofxldof assuming that the parent
// and child are of same polynomial order.
fe->GetLocalRestriction(isotr, localRVO);
// copy block
auto s = localRVO.Height() * localRVO.Width();
std::copy(localRVO.Data(), localRVO.Data() + s, bs_ptr);
bs_ptr += s;
}
DenseMatrix &lR =
fespace->IsVariableOrder() ? localRVO : localR[geom](emb.matrix);
max_rows = std::max(lR.Height(), max_rows);
auto size = lR.Height() * lR.Width();
fespace->elem_dof->GetRow(emb.parent, dofs);
if (fine_rank == MyRank)
{
// diagonal
old_elem_dof->GetRow(k, old_dofs);
MFEM_VERIFY(old_dofs.Size() == dofs.Size(),
"Parent and child must have same #dofs.");
block_row_idcs_offsets[diag_idx + 1] =
block_row_idcs_offsets[diag_idx] + lR.Height();
block_col_idcs_offsets[diag_idx + 1] =
block_col_idcs_offsets[diag_idx] + lR.Width();
if (fespace->IsVariableOrder())
{
block_offsets[diag_idx] = var_offset;
var_offset += size;
}
else
{
block_offsets[diag_idx] = geom_offsets[geom] + size * emb.matrix;
}
for (int i = 0; i < lR.Height(); ++i, ++ridx)
{
if (!std::isfinite(lR(i, 0)))
{
row_idcs[ridx] = INT_MAX;
continue;
}
int r = dofs[i];
int m = (r >= 0) ? r : (-1 - r);
if (is_dg || !mark[m])
{
row_idcs[ridx] = r;
mark[m] = 1;
}
else
{
row_idcs[ridx] = INT_MAX;
}
}
for (int i = 0; i < lR.Width(); ++i, ++cidx)
{
col_idcs[cidx] = old_dofs[i];
}
++diag_idx;
}
else
{
// off-diagonal
auto &tmp = ks_map.at(k);
auto od_idx = tmp[0];
auto od_ridx = tmp[1];
block_off_diag_row_idcs_offsets[od_idx + 1] =
block_off_diag_row_idcs_offsets[od_idx] + lR.Height();
block_off_diag_col_offsets[od_idx] = od_ridx;
block_off_diag_widths[od_idx] = lR.Width();
recv_segment_idcs[od_idx] = tmp[2];
if (fespace->IsVariableOrder())
{
off_diag_block_offsets[od_idx] = var_offset;
var_offset += size;
}
else
{
off_diag_block_offsets[od_idx] =
geom_offsets[geom] + size * emb.matrix;
}
for (int i = 0; i < lR.Height(); ++i, ++od_ridx)
{
if (!std::isfinite(lR(i, 0)))
{
row_off_diag_idcs[od_ridx] = INT_MAX;
continue;
}
int r = dofs[i];
int m = (r >= 0) ? r : (-1 - r);
if (is_dg || !mark[m])
{
row_off_diag_idcs[od_ridx] = r;
mark[m] = 1;
}
else
{
row_off_diag_idcs[od_ridx] = INT_MAX;
}
}
++od_idx;
}
}
}
// if not using GPU, set max_rows/max_cols to zero
if (Device::Allows(Backend::DEVICE_MASK))
{
max_rows = std::min(max_rows, max_team_size);
}
else
{
max_rows = 1;
}
requests.reserve(recv_ranks.Size() + send_ranks.Size());
}
} // namespace mfem
/// \endcond DO_NOT_DOCUMENT
#endif
-111
View File
@@ -1,111 +0,0 @@
// Copyright (c) 2010-2025, Lawrence Livermore National Security, LLC. Produced
// at the Lawrence Livermore National Laboratory. All Rights reserved. See files
// LICENSE and NOTICE for details. LLNL-CODE-806117.
//
// This file is part of the MFEM library. For more information and source code
// availability visit https://mfem.org.
//
// MFEM is free software; you can redistribute it and/or modify it under the
// terms of the BSD-3 license. We welcome feedback and contributions, see file
// CONTRIBUTING.md for details.
#ifndef MFEM_PDEREFMAT_OP
#define MFEM_PDEREFMAT_OP
#include "../config/config.hpp"
#ifdef MFEM_USE_MPI
#include "pfespace.hpp"
#include "kernel_dispatch.hpp"
#include <vector>
namespace mfem
{
/// \cond DO_NOT_DOCUMENT
struct ParDerefineMatrixOp : public Operator
{
ParFiniteElementSpace *fespace;
/// offsets into block_storage for diagonal
Array<int> block_offsets;
/// offsets into row_idcs for diagonal
Array<int> block_row_idcs_offsets;
/// offsets into col_idcs for diagonal
Array<int> block_col_idcs_offsets;
/// offsets into block_storage for off-diagonal
Array<int> off_diag_block_offsets;
/// offsets into row_idcs for off-diagonal
Array<int> block_off_diag_row_idcs_offsets;
Array<int> block_off_diag_col_offsets;
Array<int> block_off_diag_widths;
/// mapping for row dofs, INT_MAX indicates the block row should be ignored.
/// negative means the row data should be negated.
/// only for diagonal blocks
Array<int> row_idcs;
/// mapping for col dofs, negative means the col data should be negated.
/// only for diagonal blocks
Array<int> col_idcs;
Array<int> pack_col_idcs;
/// mapping for row dofs, INT_MAX indicates the block row should be ignored.
/// negative means the row data should be negated.
/// only for off-diagonal blocks
Array<int> row_off_diag_idcs;
/// dense block matrices which can be reused to construct the full matrix
/// operation. These are stored contiguously and blocks have no restrictions
/// on shape (can be rectangle and differ from block to block).
/// This is only for the diagonal block.
Vector block_storage;
/// maximum height of any block in block_storage for GPU
/// parallelization, or 1 for CPU runs.
int max_rows;
/// quasi Ordering::byNODES, broken into sections by ranks we need to send
/// the data to
mutable Vector xghost_send;
/// quasi Ordering::byNODES, broken into sections by ranks we received
/// the data from
mutable Vector xghost_recv;
/// maps off-diagonal k to segment
Array<int> recv_segment_idcs;
/// cumulative count of dofs which will be received from other ranks
Array<int> recv_segments;
/// Source rank of each recv segment
Array<int> recv_ranks;
/// What send segment each entry in send_permutations corresponds to
Array<int> send_segment_idcs;
/// cumulative count of dofs which will be sent to other ranks
Array<int> send_segments;
/// Destination rank of each send segment
Array<int> send_ranks;
/// how to permute/sign change values from our local x to send to other ranks
Array<int> send_permutations;
/// internal buffer for MPI requests
mutable std::vector<MPI_Request> requests;
using MultKernelType = void (*)(const ParDerefineMatrixOp &, const Vector &,
Vector &);
/// template args: ordering, atomic
MFEM_REGISTER_KERNELS(MultKernel, MultKernelType, (Ordering::Type, bool));
struct Kernels
{
Kernels();
};
void Mult(const Vector &x, Vector &y) const;
ParDerefineMatrixOp(ParFiniteElementSpace &fespace_, int old_ndofs,
const Table *old_elem_dof, const Table *old_elem_fos);
};
/// \endcond DO_NOT_DOCUMENT
} // namespace mfem
#endif
#endif
+196 -111
View File
@@ -22,13 +22,12 @@
#include "../mesh/mesh_headers.hpp"
#include "../general/binaryio.hpp"
#include "pderefmat_op.hpp"
#include <limits>
#include <list>
namespace mfem
{
ParFiniteElementSpace::ParFiniteElementSpace(
const ParFiniteElementSpace &orig, ParMesh *pmesh,
const FiniteElementCollection *fec)
@@ -4488,6 +4487,13 @@ ParFiniteElementSpace::RebalanceMatrix(int old_ndofs,
return M;
}
struct DerefDofMessage
{
std::vector<HYPRE_BigInt> dofs;
MPI_Request request;
};
HypreParMatrix*
ParFiniteElementSpace::ParallelDerefinementMatrix(int old_ndofs,
const Table* old_elem_dof,
@@ -4530,13 +4536,7 @@ ParFiniteElementSpace::ParallelDerefinementMatrix(int old_ndofs,
old_pncmesh->GetDerefinementTransforms();
const Array<int> &old_ranks = old_pncmesh->GetDerefineOldRanks();
// key: other rank
// value: send or recieve buffer
std::map<int, std::vector<HYPRE_BigInt>> to_send;
std::map<int, std::vector<HYPRE_BigInt>> to_recv;
// key: index into dtrans.embeddings
// value: [start, stop]
std::unordered_map<int, std::array<size_t, 2>> recv_messages;
std::map<int, DerefDofMessage> messages;
HYPRE_BigInt old_offset = HYPRE_AssumedPartitionCheck()
? old_dof_offsets[0] : old_dof_offsets[MyRank];
@@ -4556,46 +4556,30 @@ ParFiniteElementSpace::ParallelDerefinementMatrix(int old_ndofs,
old_elem_dof->GetRow(k, dofs);
DofsToVDofs(dofs, old_ndofs);
std::vector<HYPRE_BigInt>& send_buf = to_send[coarse_rank];
auto pos = send_buf.size();
send_buf.resize(pos + dofs.Size());
DerefDofMessage &msg = messages[k];
msg.dofs.resize(dofs.Size());
for (int i = 0; i < dofs.Size(); i++)
{
send_buf[pos + i] = old_offset + dofs[i];
msg.dofs[i] = old_offset + dofs[i];
}
MPI_Isend(&msg.dofs[0], static_cast<int>(msg.dofs.size()), HYPRE_MPI_BIG_INT,
coarse_rank, 291, MyComm, &msg.request);
}
else if (coarse_rank == MyRank && fine_rank != MyRank)
{
MFEM_ASSERT(emb.parent >= 0, "");
Geometry::Type geom = mesh->GetElementBaseGeometry(emb.parent);
std::vector<HYPRE_BigInt>& recv_buf = to_recv[fine_rank];
auto& msg = recv_messages[k];
msg[0] = recv_buf.size();
recv_buf.resize(recv_buf.size() + ldof[geom] * vdim);
msg[1] = recv_buf.size();
}
}
DerefDofMessage &msg = messages[k];
msg.dofs.resize(ldof[geom]*vdim);
// assume embedding orders are consistent (i.e. what we expect to receive
// first from a given rank is sent first, etc.)
std::vector<MPI_Request> requests;
requests.reserve(to_send.size() + to_recv.size());
// enqueue recvs
for (auto &v : to_recv)
{
requests.emplace_back();
MPI_Irecv(v.second.data(), v.second.size(), HYPRE_MPI_BIG_INT, v.first,
MessageTag::DEREFINEMENT_MATRIX_CONSTRUCTION_DATA, MyComm,
&requests.back());
}
// enqueue sends
for (auto &v : to_send)
{
requests.emplace_back();
MPI_Isend(v.second.data(), v.second.size(), HYPRE_MPI_BIG_INT, v.first,
MessageTag::DEREFINEMENT_MATRIX_CONSTRUCTION_DATA, MyComm,
&requests.back());
MPI_Irecv(&msg.dofs[0], ldof[geom]*vdim, HYPRE_MPI_BIG_INT,
fine_rank, 291, MyComm, &msg.request);
}
// TODO: coalesce Isends/Irecvs to the same rank. Typically, on uniform
// derefinement, there should be just one send to MyRank-1 and one recv
// from MyRank+1
}
DenseTensor localR[Geometry::NumGeom];
@@ -4653,7 +4637,10 @@ ParFiniteElementSpace::ParallelDerefinementMatrix(int old_ndofs,
diag->Finalize();
// wait for all sends/receives to complete
MPI_Waitall(requests.size(), requests.data(), MPI_STATUSES_IGNORE);
for (auto it = messages.begin(); it != messages.end(); ++it)
{
MPI_Wait(&it->second.request, MPI_STATUS_IGNORE);
}
// create the off-diagonal part of the derefinement matrix
SparseMatrix *offd = new SparseMatrix(ndofs*vdim, 1);
@@ -4674,14 +4661,13 @@ ParFiniteElementSpace::ParallelDerefinementMatrix(int old_ndofs,
elem_dof->GetRow(emb.parent, dofs);
auto& odofs = to_recv.at(fine_rank);
auto &msg = recv_messages[k];
MFEM_ASSERT(msg[1] > msg[0], "");
DerefDofMessage &msg = messages[k];
MFEM_ASSERT(msg.dofs.size(), "");
for (int vd = 0; vd < vdim; vd++)
{
MFEM_ASSERT(ldof[geom], "");
HYPRE_BigInt *remote_dofs = odofs.data() + msg[0] + vd * ldof[geom];
HYPRE_BigInt* remote_dofs = &msg.dofs[vd*ldof[geom]];
for (int i = 0; i < lR.Height(); i++)
{
@@ -4708,6 +4694,7 @@ ParFiniteElementSpace::ParallelDerefinementMatrix(int old_ndofs,
}
}
messages.clear();
offd->Finalize(0);
offd->SetWidth(static_cast<int>(col_map.size()));
@@ -4959,13 +4946,8 @@ void ParFiniteElementSpace::Update(bool want_transform)
case Mesh::DEREFINE:
{
#if 0
Th.Reset(ParallelDerefinementMatrix(old_ndofs, old_elem_dof,
old_elem_fos));
#else
Th.Reset(new ParDerefineMatrixOp(*this, old_ndofs, old_elem_dof,
old_elem_fos));
#endif
if (Nonconforming())
{
Th.SetOperatorOwner(false);
@@ -5327,6 +5309,18 @@ DeviceConformingProlongationOperator::DeviceConformingProlongationOperator(
if (recv_size > 0) { req_counter++; }
}
requests = new MPI_Request[req_counter];
if (mpi_gpu_aware)
{
#if defined(MFEM_USE_HIP)
MFEM_GPU_CHECK(
hipEventCreateWithFlags(&gpu_event, hipEventDisableTiming));
#elif defined(MFEM_USE_CUDA)
MFEM_GPU_CHECK(
cudaEventCreateWithFlags(&gpu_event, cudaEventDisableTiming));
#else
MFEM_ABORT("not implemented");
#endif
}
}
DeviceConformingProlongationOperator::DeviceConformingProlongationOperator(
@@ -5340,16 +5334,27 @@ DeviceConformingProlongationOperator::DeviceConformingProlongationOperator(
}
static void ExtractSubVector(const Array<int> &indices,
const Vector &vin, Vector &vout)
const Vector &vin, Vector &vout,
real_t a, real_t b)
{
MFEM_ASSERT(indices.Size() == vout.Size(), "incompatible sizes!");
auto y = vout.Write();
auto y = (a == 0) ? vout.Write() : vout.ReadWrite();
const auto x = vin.Read();
const auto I = indices.Read();
mfem::forall(indices.Size(), [=] MFEM_HOST_DEVICE (int i)
if (a == 0)
{
y[i] = x[I[i]];
}); // indices can be repeated
mfem::forall(indices.Size(), [=] MFEM_HOST_DEVICE (int i)
{
y[i] = b*x[I[i]];
}); // indices can be repeated
}
else
{
mfem::forall(indices.Size(), [=] MFEM_HOST_DEVICE (int i)
{
y[i] = a*y[i] + b*x[I[i]];
}); // indices can be repeated
}
}
void DeviceConformingProlongationOperator::BcastBeginCopy(
@@ -5357,10 +5362,7 @@ void DeviceConformingProlongationOperator::BcastBeginCopy(
{
// shr_buf[i] = src[shr_ltdof[i]]
if (shr_ltdof.Size() == 0) { return; }
ExtractSubVector(shr_ltdof, x, shr_buf);
// If the above kernel is executed asynchronously, we should wait for it to
// complete
if (mpi_gpu_aware) { MFEM_STREAM_SYNC; }
ExtractSubVector(shr_ltdof, x, shr_buf, 0, 1);
}
static void SetSubVector(const Array<int> &indices,
@@ -5397,7 +5399,7 @@ void DeviceConformingProlongationOperator::Mult(const Vector &x,
Vector &y) const
{
const GroupTopology &gtopo = gc.GetGroupTopology();
int req_counter = 0;
int req_counter = 0, num_recv_req = 0;
// Make sure 'y' is marked as valid on device and for use on device.
// This ensures that there is no unnecessary host to device copy when the
// input 'y' is valid on host (in 'y.SetSubVector(ext_ldof, 0.0)' when local
@@ -5407,42 +5409,89 @@ void DeviceConformingProlongationOperator::Mult(const Vector &x,
{
// done on device since we've marked ext_ldof for use on device:
y.SetSubVector(ext_ldof, 0.0);
BcastLocalCopy(x, y);
return;
}
else
BcastBeginCopy(x); // copy to 'shr_buf'
if (mpi_gpu_aware && shr_ltdof.Size() != 0)
{
BcastBeginCopy(x); // copy to 'shr_buf'
/* record a stream event to wait for later */
#if defined(MFEM_USE_HIP)
MFEM_GPU_CHECK(hipEventRecord(gpu_event, 0));
#elif defined(MFEM_USE_CUDA)
MFEM_GPU_CHECK(cudaEventRecord(gpu_event, 0));
#endif
}
BcastLocalCopy(x, y);
// Queue all receive communications
if (ext_ldof.Size() != 0) // ext_ldof.Size() == ext_buf.Size()
{
auto recv_buf = mpi_gpu_aware ? ext_buf.Write() : ext_buf.HostWrite();
for (int nbr = 1; nbr < gtopo.GetNumNeighbors(); nbr++)
{
const int recv_offset = ext_buf_offsets[nbr];
const int recv_size = ext_buf_offsets[nbr+1] - recv_offset;
if (recv_size > 0)
{
MPI_Irecv(recv_buf + recv_offset, recv_size,
MPITypeMap<real_t>::mpi_type,
gtopo.GetNeighborRank(nbr), 41822,
gtopo.GetComm(), &requests[req_counter++]);
}
}
num_recv_req = req_counter;
}
// Queue all send communications
if (shr_ltdof.Size() != 0) // shr_ltdof.Size() == shr_buf.Size()
{
// The BcastBeginCopy kernel is executed asynchronously, we should wait
// for it to complete:
// - when mpi_gpu_aware == false, this is done implicily when we call
// shr_buf.HostRead()
// - when mpi_gpu_aware == true, we need to wait for BcastBeginCopy to
// complete by waiting for gpu_event.
if (mpi_gpu_aware)
{
/* wait for the stream event recorded above */
#if defined(MFEM_USE_HIP)
MFEM_GPU_CHECK(hipEventSynchronize(gpu_event));
#elif defined(MFEM_USE_CUDA)
MFEM_GPU_CHECK(cudaEventSynchronize(gpu_event));
#endif
}
auto send_buf = mpi_gpu_aware ? shr_buf.Read() : shr_buf.HostRead();
for (int nbr = 1; nbr < gtopo.GetNumNeighbors(); nbr++)
{
const int send_offset = shr_buf_offsets[nbr];
const int send_size = shr_buf_offsets[nbr+1] - send_offset;
if (send_size > 0)
{
auto send_buf = mpi_gpu_aware ? shr_buf.Read() : shr_buf.HostRead();
MPI_Isend(send_buf + send_offset, send_size, MPITypeMap<real_t>::mpi_type,
gtopo.GetNeighborRank(nbr), 41822,
gtopo.GetComm(), &requests[req_counter++]);
}
const int recv_offset = ext_buf_offsets[nbr];
const int recv_size = ext_buf_offsets[nbr+1] - recv_offset;
if (recv_size > 0)
{
auto recv_buf = mpi_gpu_aware ? ext_buf.Write() : ext_buf.HostWrite();
MPI_Irecv(recv_buf + recv_offset, recv_size, MPITypeMap<real_t>::mpi_type,
MPI_Isend(send_buf + send_offset, send_size,
MPITypeMap<real_t>::mpi_type,
gtopo.GetNeighborRank(nbr), 41822,
gtopo.GetComm(), &requests[req_counter++]);
}
}
}
BcastLocalCopy(x, y);
if (!local)
{
MPI_Waitall(req_counter, requests, MPI_STATUSES_IGNORE);
BcastEndCopy(y); // copy from 'ext_buf'
}
// Wait for all receive requests
MPI_Waitall(num_recv_req, requests, MPI_STATUSES_IGNORE);
BcastEndCopy(y); // copy from 'ext_buf'
// Wait for all send requests
MPI_Waitall(req_counter - num_recv_req, requests + num_recv_req,
MPI_STATUSES_IGNORE);
}
DeviceConformingProlongationOperator::~DeviceConformingProlongationOperator()
{
if (mpi_gpu_aware)
{
#if defined(MFEM_USE_HIP)
MFEM_GPU_CHECK(hipEventDestroy(gpu_event));
#elif defined(MFEM_USE_CUDA)
MFEM_GPU_CHECK(cudaEventDestroy(gpu_event));
#endif
}
delete [] requests;
ext_buf_offsets.Delete();
shr_buf_offsets.Delete();
@@ -5453,25 +5502,23 @@ void DeviceConformingProlongationOperator::ReduceBeginCopy(
{
// ext_buf[i] = src[ext_ldof[i]]
if (ext_ldof.Size() == 0) { return; }
ExtractSubVector(ext_ldof, x, ext_buf);
// If the above kernel is executed asynchronously, we should wait for it to
// complete
if (mpi_gpu_aware) { MFEM_STREAM_SYNC; }
ExtractSubVector(ext_ldof, x, ext_buf, 0, 1);
}
void DeviceConformingProlongationOperator::ReduceLocalCopy(
const Vector &x, Vector &y) const
const Vector &x, Vector &y, real_t a, real_t b) const
{
// dst[i] = src[ltdof_ldof[i]]
if (ltdof_ldof.Size() == 0) { return; }
ExtractSubVector(ltdof_ldof, x, y);
ExtractSubVector(ltdof_ldof, x, y, a, b);
}
static void AddSubVector(const Array<int> &unique_dst_indices,
const Array<int> &unique_to_src_offsets,
const Array<int> &unique_to_src_indices,
const Vector &src,
Vector &dst)
Vector &dst,
real_t b)
{
auto y = dst.ReadWrite();
const auto x = src.Read();
@@ -5481,56 +5528,94 @@ static void AddSubVector(const Array<int> &unique_dst_indices,
mfem::forall(unique_dst_indices.Size(), [=] MFEM_HOST_DEVICE (int i)
{
const int dst_idx = DST_I[i];
real_t sum = y[dst_idx];
real_t sum = 0;
const int end = SRC_O[i+1];
for (int j = SRC_O[i]; j != end; ++j) { sum += x[SRC_I[j]]; }
y[dst_idx] = sum;
y[dst_idx] += b*sum;
});
}
void DeviceConformingProlongationOperator::ReduceEndAssemble(Vector &y) const
void DeviceConformingProlongationOperator::ReduceEndAssemble(
Vector &y, real_t b) const
{
// dst[shr_ltdof[i]] += shr_buf[i]
if (unq_ltdof.Size() == 0) { return; }
AddSubVector(unq_ltdof, unq_shr_i, unq_shr_j, shr_buf, y);
AddSubVector(unq_ltdof, unq_shr_i, unq_shr_j, shr_buf, y, b);
}
void DeviceConformingProlongationOperator::MultTranspose(const Vector &x,
Vector &y) const
void DeviceConformingProlongationOperator::ApplyTranspose(
const Vector &x, Vector &y, real_t a, real_t b) const
{
const GroupTopology &gtopo = gc.GetGroupTopology();
int req_counter = 0;
if (!local)
int req_counter = 0, num_recv_req = 0;
if (local)
{
ReduceBeginCopy(x); // copy to 'ext_buf'
ReduceLocalCopy(x, y, a, b);
return;
}
ReduceBeginCopy(x); // copy to 'ext_buf'
if (mpi_gpu_aware && ext_ldof.Size() != 0)
{
/* record a stream event to wait for later */
#if defined(MFEM_USE_HIP)
MFEM_GPU_CHECK(hipEventRecord(gpu_event, 0));
#elif defined(MFEM_USE_CUDA)
MFEM_GPU_CHECK(cudaEventRecord(gpu_event, 0));
#endif
}
ReduceLocalCopy(x, y, a, b);
// Queue all receive communications
if (unq_ltdof.Size() != 0)
{
auto recv_buf = mpi_gpu_aware ? shr_buf.Write() : shr_buf.HostWrite();
for (int nbr = 1; nbr < gtopo.GetNumNeighbors(); nbr++)
{
const int recv_offset = shr_buf_offsets[nbr];
const int recv_size = shr_buf_offsets[nbr+1] - recv_offset;
if (recv_size > 0)
{
MPI_Irecv(recv_buf + recv_offset, recv_size,
MPITypeMap<real_t>::mpi_type,
gtopo.GetNeighborRank(nbr), 41823,
gtopo.GetComm(), &requests[req_counter++]);
}
}
num_recv_req = req_counter;
}
// Queue all send communications
if (ext_ldof.Size() != 0)
{
if (mpi_gpu_aware)
{
/* wait for the stream event recorded above */
#if defined(MFEM_USE_HIP)
MFEM_GPU_CHECK(hipEventSynchronize(gpu_event));
#elif defined(MFEM_USE_CUDA)
MFEM_GPU_CHECK(cudaEventSynchronize(gpu_event));
#endif
}
auto send_buf = mpi_gpu_aware ? ext_buf.Read() : ext_buf.HostRead();
for (int nbr = 1; nbr < gtopo.GetNumNeighbors(); nbr++)
{
const int send_offset = ext_buf_offsets[nbr];
const int send_size = ext_buf_offsets[nbr+1] - send_offset;
if (send_size > 0)
{
auto send_buf = mpi_gpu_aware ? ext_buf.Read() : ext_buf.HostRead();
MPI_Isend(send_buf + send_offset, send_size, MPITypeMap<real_t>::mpi_type,
gtopo.GetNeighborRank(nbr), 41823,
gtopo.GetComm(), &requests[req_counter++]);
}
const int recv_offset = shr_buf_offsets[nbr];
const int recv_size = shr_buf_offsets[nbr+1] - recv_offset;
if (recv_size > 0)
{
auto recv_buf = mpi_gpu_aware ? shr_buf.Write() : shr_buf.HostWrite();
MPI_Irecv(recv_buf + recv_offset, recv_size, MPITypeMap<real_t>::mpi_type,
MPI_Isend(send_buf + send_offset, send_size,
MPITypeMap<real_t>::mpi_type,
gtopo.GetNeighborRank(nbr), 41823,
gtopo.GetComm(), &requests[req_counter++]);
}
}
}
ReduceLocalCopy(x, y);
if (!local)
{
MPI_Waitall(req_counter, requests, MPI_STATUSES_IGNORE);
ReduceEndAssemble(y); // assemble from 'shr_buf'
}
// Wait for all receive requests
MPI_Waitall(num_recv_req, requests, MPI_STATUSES_IGNORE);
ReduceEndAssemble(y, b); // assemble from 'shr_buf'
// Wait for all send requests
MPI_Waitall(req_counter - num_recv_req, requests + num_recv_req,
MPI_STATUSES_IGNORE);
}
} // namespace mfem
+18 -6
View File
@@ -24,12 +24,9 @@
namespace mfem
{
struct ParDerefineMatrixOp;
/// Abstract parallel finite element space.
class ParFiniteElementSpace : public FiniteElementSpace
{
friend struct ParDerefineMatrixOp;
private:
/// MPI data.
MPI_Comm MyComm;
@@ -589,6 +586,7 @@ public:
{ MultTranspose(x,y); }
};
/// Auxiliary device class used by ParFiniteElementSpace.
class DeviceConformingProlongationOperator: public
ConformingProlongationOperator
@@ -601,6 +599,11 @@ protected:
Array<int> ltdof_ldof, unq_ltdof;
Array<int> unq_shr_i, unq_shr_j;
MPI_Request *requests;
#if defined(MFEM_USE_HIP)
hipEvent_t gpu_event;
#elif defined(MFEM_USE_CUDA)
cudaEvent_t gpu_event;
#endif
// Kernel: copy ltdofs from 'src' to 'shr_buf' - prepare for send.
// shr_buf[i] = src[shr_ltdof[i]]
@@ -620,11 +623,12 @@ protected:
// Kernel: copy owned ldofs from 'src' to ltdofs in 'dst'.
// dst[i] = src[ltdof_ldof[i]]
void ReduceLocalCopy(const Vector &src, Vector &dst) const;
void ReduceLocalCopy(const Vector &src, Vector &dst,
real_t a, real_t b) const;
// Kernel: assemble dofs from 'shr_buf' into to 'dst' - after recv.
// dst[shr_ltdof[i]] += shr_buf[i]
void ReduceEndAssemble(Vector &dst) const;
void ReduceEndAssemble(Vector &dst, real_t b) const;
public:
DeviceConformingProlongationOperator(
@@ -635,12 +639,20 @@ public:
virtual ~DeviceConformingProlongationOperator();
void ApplyTranspose(const Vector &x, Vector &y,
real_t a, real_t b) const;
void Mult(const Vector &x, Vector &y) const override;
void AbsMult(const Vector &x, Vector &y) const override
{ Mult(x,y); }
void MultTranspose(const Vector &x, Vector &y) const override;
void MultTranspose(const Vector &x, Vector &y) const override
{ ApplyTranspose(x, y, 0, 1); }
void AddMultTranspose(const Vector &x, Vector &y,
real_t a = 1) const override
{ ApplyTranspose(x, y, 1, a); }
void AbsMultTranspose(const Vector &x, Vector &y) const override
{ MultTranspose(x,y); }
+10 -10
View File
@@ -214,16 +214,16 @@ public:
inline void QuadratureFunction::GetValues(
int idx, Vector &values)
{
const int s_offset = qspace->Offset(idx);
const int sl_size = qspace->Offset(idx + 1) - s_offset;
const int s_offset = qspace->offsets[idx];
const int sl_size = qspace->offsets[idx+1] - s_offset;
values.MakeRef(*this, vdim*s_offset, vdim*sl_size);
}
inline void QuadratureFunction::GetValues(
int idx, Vector &values) const
{
const int s_offset = qspace->Offset(idx);
const int sl_size = qspace->Offset(idx + 1) - s_offset;
const int s_offset = qspace->offsets[idx];
const int sl_size = qspace->offsets[idx+1] - s_offset;
values.SetSize(vdim*sl_size);
values.HostWrite();
const real_t *q = HostRead() + vdim*s_offset;
@@ -236,14 +236,14 @@ inline void QuadratureFunction::GetValues(
inline void QuadratureFunction::GetValues(
int idx, const int ip_num, Vector &values)
{
const int s_offset = qspace->Offset(idx) * vdim + ip_num * vdim;
const int s_offset = qspace->offsets[idx] * vdim + ip_num * vdim;
values.MakeRef(*this, s_offset, vdim);
}
inline void QuadratureFunction::GetValues(
int idx, const int ip_num, Vector &values) const
{
const int s_offset = qspace->Offset(idx) * vdim + ip_num * vdim;
const int s_offset = qspace->offsets[idx] * vdim + ip_num * vdim;
values.SetSize(vdim);
values.HostWrite();
const real_t *q = HostRead() + s_offset;
@@ -256,8 +256,8 @@ inline void QuadratureFunction::GetValues(
inline void QuadratureFunction::GetValues(
int idx, DenseMatrix &values)
{
const int s_offset = qspace->Offset(idx);
const int sl_size = qspace->Offset(idx + 1) - s_offset;
const int s_offset = qspace->offsets[idx];
const int sl_size = qspace->offsets[idx+1] - s_offset;
// Make the values matrix memory an alias of the quadrature function memory
Memory<real_t> &values_mem = values.GetMemory();
values_mem.Delete();
@@ -268,8 +268,8 @@ inline void QuadratureFunction::GetValues(
inline void QuadratureFunction::GetValues(
int idx, DenseMatrix &values) const
{
const int s_offset = qspace->Offset(idx);
const int sl_size = qspace->Offset(idx + 1) - s_offset;
const int s_offset = qspace->offsets[idx];
const int sl_size = qspace->offsets[idx+1] - s_offset;
values.SetSize(vdim, sl_size);
values.HostWrite();
const real_t *q = HostRead() + vdim*s_offset;
+56 -75
View File
@@ -17,9 +17,8 @@ namespace mfem
{
QuadratureSpaceBase::QuadratureSpaceBase(Mesh &mesh_, Geometry::Type geom,
const IntegrationRule &ir,
QSpaceStorage storage)
: mesh(mesh_), order(ir.GetOrder()), storage(storage)
const IntegrationRule &ir)
: mesh(mesh_), order(ir.GetOrder())
{
for (int g = 0; g < Geometry::NumGeom; g++)
{
@@ -97,37 +96,27 @@ void QuadratureSpaceBase::Integrate(VectorCoefficient &coeff,
void QuadratureSpace::ConstructOffsets()
{
const int num_elem = ne;
if (storage == QSpaceStorage::COMPRESSED &&
mesh.GetNumGeometries(mesh.Dimension()) == 1)
const int num_elem = mesh.GetNE();
offsets.SetSize(num_elem + 1);
int offset = 0;
for (int i = 0; i < num_elem; i++)
{
Array<Geometry::Type> geoms;
mesh.GetGeometries(mesh.Dimension(), geoms);
offsets.SetSize(1);
offsets.HostWrite();
offsets[0] = int_rule[geoms[0]]->GetNPoints();
size = num_elem * offsets[0];
}
else
{
offsets.SetSize(num_elem + 1);
int offset = 0;
for (int i = 0; i < num_elem; i++)
{
offsets[i] = offset;
const Geometry::Type geom = mesh.GetElementBaseGeometry(i);
MFEM_ASSERT(int_rule[geom] != NULL, "Missing integration rule.");
offset += int_rule[geom]->GetNPoints();
}
offsets[num_elem] = offset;
size = offsets.Last();
offsets[i] = offset;
int geom = mesh.GetElementBaseGeometry(i);
MFEM_ASSERT(int_rule[geom] != NULL, "Missing integration rule.");
offset += int_rule[geom]->GetNPoints();
}
offsets[num_elem] = size = offset;
}
QuadratureSpace::QuadratureSpace(Mesh *mesh_, std::istream &in,
QSpaceStorage storage)
: QuadratureSpaceBase(*mesh_, 0, storage)
void QuadratureSpace::Construct()
{
ConstructIntRules(mesh.Dimension());
ConstructOffsets();
}
QuadratureSpace::QuadratureSpace(Mesh *mesh_, std::istream &in)
: QuadratureSpaceBase(*mesh_)
{
const char *msg = "invalid input stream";
std::string ident;
@@ -146,24 +135,15 @@ QuadratureSpace::QuadratureSpace(Mesh *mesh_, std::istream &in,
return;
}
ne = mesh.GetNE();
ConstructIntRules(mesh.Dimension());
Construct();
}
QuadratureSpace::QuadratureSpace(Mesh *mesh_, int order_, QSpaceStorage storage)
: QuadratureSpaceBase(*mesh_, order_, storage)
{
ne = mesh.GetNE();
ConstructIntRules(mesh.Dimension());
}
QuadratureSpace::QuadratureSpace(Mesh &mesh_, const IntegrationRule &ir,
QSpaceStorage storage)
: QuadratureSpaceBase(mesh_, mesh_.GetTypicalElementGeometry(), ir, storage)
QuadratureSpace::QuadratureSpace(Mesh &mesh_, const IntegrationRule &ir)
: QuadratureSpaceBase(mesh_, mesh_.GetTypicalElementGeometry(), ir)
{
MFEM_VERIFY(mesh.GetNumGeometries(mesh.Dimension()) <= 1,
"Constructor not valid for mixed meshes");
ne = mesh.GetNE();
ConstructOffsets();
}
void QuadratureSpace::Save(std::ostream &os) const
@@ -185,53 +165,54 @@ const Vector &QuadratureSpace::GetGeometricFactorWeights() const
}
FaceQuadratureSpace::FaceQuadratureSpace(Mesh &mesh_, int order_,
FaceType face_type_,
QSpaceStorage storage)
: QuadratureSpaceBase(mesh_, order_, storage), face_type(face_type_),
face_indices(mesh.GetFaceIndices(face_type_)),
face_indices_inv(mesh.GetInvFaceIndices(face_type_))
FaceType face_type_)
: QuadratureSpaceBase(mesh_, order_),
face_type(face_type_),
num_faces(mesh.GetNFbyType(face_type))
{
ne = face_indices.Size();
ConstructIntRules(mesh.Dimension() - 1);
Construct();
}
FaceQuadratureSpace::FaceQuadratureSpace(Mesh &mesh_, const IntegrationRule &ir,
FaceType face_type_,
QSpaceStorage storage)
: QuadratureSpaceBase(mesh_, mesh_.GetTypicalFaceGeometry(), ir, storage),
face_type(face_type_), face_indices(mesh.GetFaceIndices(face_type_)),
face_indices_inv(mesh.GetInvFaceIndices(face_type_))
FaceType face_type_)
: QuadratureSpaceBase(mesh_, mesh_.GetTypicalFaceGeometry(), ir),
face_type(face_type_),
num_faces(mesh.GetNFbyType(face_type))
{
MFEM_VERIFY(mesh.GetNumGeometries(mesh.Dimension() - 1) <= 1,
"Constructor not valid for mixed meshes");
ne = face_indices.Size();
ConstructOffsets();
}
void FaceQuadratureSpace::ConstructOffsets()
{
if (storage == QSpaceStorage::COMPRESSED &&
mesh.GetNumGeometries(mesh.Dimension() - 1) == 1)
face_indices.SetSize(num_faces);
offsets.SetSize(num_faces + 1);
int offset = 0;
int f_idx = 0;
for (int i = 0; i < mesh.GetNumFacesWithGhost(); i++)
{
Array<Geometry::Type> geoms;
mesh.GetGeometries(mesh.Dimension() - 1, geoms);
offsets.SetSize(1);
offsets.HostWrite();
offsets[0] = int_rule[geoms[0]]->GetNPoints();
size = ne * offsets[0];
}
else
{
offsets.SetSize(face_indices.Size() + 1);
int offset = 0;
for (int i = 0; i < mesh.GetNFbyType(face_type); ++i)
const Mesh::FaceInformation face = mesh.GetFaceInformation(i);
if (face.IsNonconformingCoarse() || !face.IsOfFaceType(face_type))
{
offsets[i] = offset;
Geometry::Type geom = mesh.GetFaceGeometry(face_indices[i]);
MFEM_ASSERT(int_rule[geom] != NULL, "Missing integration rule");
offset += int_rule[geom]->GetNPoints();
continue;
}
offsets[face_indices.Size()] = size = offset;
face_indices[f_idx] = i;
face_indices_inv[i] = f_idx;
offsets[f_idx] = offset;
Geometry::Type geom = mesh.GetFaceGeometry(i);
MFEM_ASSERT(int_rule[geom] != NULL, "Missing integration rule");
offset += int_rule[geom]->GetNPoints();
f_idx++;
}
offsets[num_faces] = size = offset;
}
void FaceQuadratureSpace::Construct()
{
ConstructIntRules(mesh.Dimension() - 1);
ConstructOffsets();
}
int FaceQuadratureSpace::GetPermutedIndex(int idx, int iq) const
+25 -80
View File
@@ -19,49 +19,35 @@
namespace mfem
{
enum class QSpaceStorage
{
FULL,
COMPRESSED
};
/// Abstract base class for QuadratureSpace and FaceQuadratureSpace.
/** This class represents the storage layout for QuadratureFunction%s, that may
be defined either on mesh elements or mesh faces. */
class QuadratureSpaceBase
{
protected:
friend class QuadratureFunction; // Uses the offsets.
Mesh &mesh; ///< The underlying mesh.
int order; ///< The order of integration rule.
int size = -1; ///< Total number of quadrature points. -1 indicates
///< offsets/size not computed yet.
int ne; ///< Actual number of entities
int size; ///< Total number of quadrature points.
mutable Vector weights; ///< Integration weights.
mutable long nodes_sequence = 0; ///< Nodes counter for cache invalidation.
QSpaceStorage storage;
/// @brief Entity quadrature point offset array.
///
/// Supports a constant compression scheme for meshes which have a single
/// geometry type. When compressed, will have a single value. The true offset
/// can be computed as i * offsets[0], where i is the entity index. Otherwise
/// has size num_entities + 1. Lazily constructed.
/// @brief Entity quadrature point offset array, of size num_entities + 1.
///
/// The quadrature point values for entity i are stored in the indices between
/// offsets[i] and offsets[i+1].
Array<int> offsets;
/// The quadrature rules used for each geometry type.
const IntegrationRule *int_rule[Geometry::NumGeom];
/// Protected constructor. Used by derived classes.
QuadratureSpaceBase(Mesh &mesh_, int order_ = 0,
QSpaceStorage storage = QSpaceStorage::COMPRESSED)
: mesh(mesh_), order(order_), storage(storage)
{}
QuadratureSpaceBase(Mesh &mesh_, int order_ = 0)
: mesh(mesh_), order(order_) { }
/// Protected constructor. Used by derived classes.
QuadratureSpaceBase(Mesh &mesh_, Geometry::Type geom,
const IntegrationRule &ir,
QSpaceStorage storage = QSpaceStorage::COMPRESSED);
const IntegrationRule &ir);
/// Fill the @ref int_rule array for each geometry type using @ref order.
void ConstructIntRules(int dim);
@@ -72,55 +58,15 @@ protected:
/// Compute the integration weights.
void ConstructWeights() const;
virtual void ConstructOffsets() = 0;
public:
QSpaceStorage StorageType() const { return storage; }
/// @brief Gets the offset for a given entity @a idx.
///
/// The quadrature point values for entity i are stored in the indices
/// between Offset(i) and Offset(i+1)
int Offset(int idx) const
{
if (size < 0)
{
const_cast<QuadratureSpaceBase *>(this)->ConstructOffsets();
}
return (offsets.Size() == 1) ? (idx * offsets[0]) : offsets[idx];
}
/// @brief Entity quadrature point offset array.
///
/// Supports a constant compression scheme for meshes which have a single
/// geometry type. When compressed, will have a single value. The true offset
/// can be computed as i * offsets[0], where i is the entity index. Otherwise
/// has size num_entities + 1.
///
const Array<int> &Offsets() const
{
if (size < 0)
{
const_cast<QuadratureSpaceBase *>(this)->ConstructOffsets();
}
return offsets;
}
/// Return the total number of quadrature points.
int GetSize() const
{
if (size < 0)
{
const_cast<QuadratureSpaceBase *>(this)->ConstructOffsets();
}
return size;
}
int GetSize() const { return size; }
/// Return the order of the quadrature rule(s) used by all elements.
int GetOrder() const { return order; }
/// Return the number of entities.
int GetNE() const { return ne; }
int GetNE() const { return offsets.Size() - 1; }
/// Returns the mesh.
inline Mesh *GetMesh() const { return &mesh; }
@@ -174,20 +120,19 @@ class QuadratureSpace : public QuadratureSpaceBase
{
protected:
const Vector &GetGeometricFactorWeights() const override;
void ConstructOffsets() override;
void ConstructOffsets();
void Construct();
public:
/// Create a QuadratureSpace based on the global rules from #IntRules.
QuadratureSpace(Mesh *mesh_, int order_,
QSpaceStorage storage = QSpaceStorage::COMPRESSED);
QuadratureSpace(Mesh *mesh_, int order_)
: QuadratureSpaceBase(*mesh_, order_) { Construct(); }
/// @brief Create a QuadratureSpace with an IntegrationRule, valid only when
/// the mesh has one element type.
QuadratureSpace(Mesh &mesh_, const IntegrationRule &ir,
QSpaceStorage storage = QSpaceStorage::COMPRESSED);
QuadratureSpace(Mesh &mesh_, const IntegrationRule &ir);
/// Read a QuadratureSpace from the stream @a in.
QuadratureSpace(Mesh *mesh_, std::istream &in,
QSpaceStorage storage = QSpaceStorage::COMPRESSED);
QuadratureSpace(Mesh *mesh_, std::istream &in);
/// Returns number of elements in the mesh.
inline int GetNE() const { return mesh.GetNE(); }
@@ -224,29 +169,29 @@ public:
class FaceQuadratureSpace : public QuadratureSpaceBase
{
FaceType face_type; ///< Is the space defined on interior or boundary faces?
const int num_faces; ///< Number of faces.
/// Map from boundary or interior face indices to mesh face indices.
const Array<int> &face_indices;
Array<int> face_indices;
/// Inverse of the map @a face_indices.
const std::unordered_map<int,int> &face_indices_inv;
std::unordered_map<int,int> face_indices_inv;
const Vector &GetGeometricFactorWeights() const override;
void ConstructOffsets() override;
void ConstructOffsets();
void Construct();
public:
/// Create a FaceQuadratureSpace based on the global rules from #IntRules.
FaceQuadratureSpace(Mesh &mesh_, int order_, FaceType face_type_,
QSpaceStorage storage = QSpaceStorage::COMPRESSED);
FaceQuadratureSpace(Mesh &mesh_, int order_, FaceType face_type_);
/// @brief Create a FaceQuadratureSpace with an IntegrationRule, valid only
/// when the mesh has one type of face geometry.
FaceQuadratureSpace(Mesh &mesh_, const IntegrationRule &ir,
FaceType face_type_,
QSpaceStorage storage = QSpaceStorage::COMPRESSED);
FaceType face_type_);
/// Returns number of faces in the mesh.
inline int GetNumFaces() const { return face_indices.Size(); }
inline int GetNumFaces() const { return num_faces; }
/// Returns the face type (boundary or interior).
FaceType GetFaceType() const { return face_type; }
+2
View File
@@ -503,6 +503,8 @@ void QuadratureInterpolator::Mult(const Vector &e_vec,
Vector &q_der,
Vector &q_det) const
{
MFEM_PERF_FUNCTION;
using namespace internal::quadrature_interpolator;
const int ne = fespace->GetNE();
+181 -10
View File
@@ -13,6 +13,7 @@
#include "bilinearform.hpp"
#include "pbilinearform.hpp"
#include "../general/forall.hpp"
#include "kernels.hpp"
namespace mfem
{
@@ -2322,6 +2323,74 @@ void Prolongation2D(const int NE, const int D1D, const int Q1D,
});
}
template <int DLO, int DHI>
static void SmemProlongation3D(const int NE,
const Vector& localL, Vector& localH,
const Array<real_t> &b, const Vector& mask)
{
auto u_lo = Reshape(localL.Read(), DLO, DLO, DLO, NE);
auto u_hi = Reshape(localH.Write(), DHI, DHI, DHI, NE);
auto d_b = Reshape(b.Read(), DHI, DLO);
auto m_ = Reshape(mask.Read(), DHI, DHI, DHI, NE);
mfem::forall_2D(NE, DHI, DHI, [=] MFEM_HOST_DEVICE (int e)
{
// Load B into shared memory
MFEM_SHARED real_t s_B[DHI*DLO];
kernels::internal::LoadBt<DLO,DHI>(DLO,DHI,d_b,s_B);
const DeviceMatrix B(s_B, DHI, DLO);
MFEM_SHARED real_t s_u[DHI*DHI*DLO];
const DeviceCube u(s_u, DHI, DHI, DLO);
real_t v[DHI];
MFEM_FOREACH_THREAD(lx,x,DLO)
{
MFEM_FOREACH_THREAD(ly,y,DLO)
{
for (int hz = 0; hz < DHI; ++hz) { v[hz] = 0.0; }
for (int lz = 0; lz < DLO; ++lz)
{
const real_t XYZ = u_lo(lx,ly,lz,e);
for (int hz = 0; hz < DHI; ++hz) { v[hz] += XYZ * B(hz,lz); }
}
for (int hz = 0; hz < DHI; ++hz) { u(hz,ly,lx) = v[hz]; }
}
}
MFEM_SYNC_THREAD;
MFEM_FOREACH_THREAD(hz,y,DHI)
{
MFEM_FOREACH_THREAD(lx,x,DLO)
{
for (int hy = 0; hy < DHI; ++hy) { v[hy] = 0.0; }
for (int ly = 0; ly < DLO; ++ly)
{
const real_t zYX = u(hz,ly,lx);
for (int hy = 0; hy < DHI; ++hy) { v[hy] += zYX * B(hy,ly); }
}
for (int hy = 0; hy < DHI; ++hy) { u(hz,hy,lx) = v[hy]; }
}
}
MFEM_SYNC_THREAD;
MFEM_FOREACH_THREAD(hz,y,DHI)
{
MFEM_FOREACH_THREAD(hy,x,DHI)
{
for (int hx = 0; hx < DHI; ++hx) { v[hx] = 0.0; }
for (int lx = 0; lx < DLO; ++lx)
{
const real_t zyX = u(hz,hy,lx);
for (int hx = 0; hx < DHI; ++hx) { v[hx] += zyX * B(hx,lx); }
}
for (int hx = 0; hx < DHI; ++hx)
{
u_hi(hx,hy,hz,e) = m_(hx,hy,hz,e)*v[hx];
}
}
}
});
}
void Prolongation3D(const int NE, const int D1D, const int Q1D,
const Vector& localL, Vector& localH,
const Array<real_t>& B, const Vector& mask)
@@ -2403,9 +2472,9 @@ void Prolongation3D(const int NE, const int D1D, const int Q1D,
});
}
void Restriction2D(const int NE, const int D1D, const int Q1D,
const Vector& localH, Vector& localL,
const Array<real_t>& Bt, const Vector& mask)
void ProlongationTranspose2D(const int NE, const int D1D, const int Q1D,
const Vector& localH, Vector& localL,
const Array<real_t>& Bt, const Vector& mask)
{
auto x_ = Reshape(localH.Read(), Q1D, Q1D, NE);
auto y_ = Reshape(localL.Write(), D1D, D1D, NE);
@@ -2448,9 +2517,78 @@ void Restriction2D(const int NE, const int D1D, const int Q1D,
}
});
}
void Restriction3D(const int NE, const int D1D, const int Q1D,
const Vector& localH, Vector& localL,
const Array<real_t>& Bt, const Vector& mask)
template <int DLO, int DHI>
static void SmemProlongationTranspose3D(
const int NE, const Vector& localH, Vector& localL,
const Array<real_t>& bt, const Vector& mask)
{
auto u_h = Reshape(localH.Read(), DHI, DHI, DHI, NE);
auto u_l = Reshape(localL.Write(), DLO, DLO, DLO, NE);
auto d_bt = Reshape(bt.Read(), DLO, DHI);
auto m_ = Reshape(mask.Read(), DHI, DHI, DHI, NE);
mfem::forall_2D(NE, DHI, DHI, [=] MFEM_HOST_DEVICE (int e)
{
// Load Bt into shared memory
MFEM_SHARED real_t s_Bt[DHI*DLO];
kernels::internal::LoadBt<DHI,DLO>(DHI,DLO,d_bt,s_Bt);
const DeviceMatrix Bt(s_Bt, DLO, DHI);
MFEM_SHARED real_t s_u[DLO*DHI*DHI];
const DeviceCube u(s_u, DLO, DHI, DHI);
real_t v[DLO];
MFEM_FOREACH_THREAD(hx,x,DHI)
{
MFEM_FOREACH_THREAD(hy,y,DHI)
{
for (int lz = 0; lz < DLO; ++lz) { v[lz] = 0.0; }
for (int hz = 0; hz < DHI; ++hz)
{
const real_t XYZ = m_(hx,hy,hz,e)*u_h(hx,hy,hz,e);
for (int lz = 0; lz < DLO; ++lz) { v[lz] += XYZ * Bt(lz,hz); }
}
for (int lz = 0; lz < DLO; ++lz) { u(lz,hy,hx) = v[lz]; }
}
}
MFEM_SYNC_THREAD;
MFEM_FOREACH_THREAD(lz,y,DLO)
{
MFEM_FOREACH_THREAD(hx,x,DHI)
{
for (int ly = 0; ly < DLO; ++ly) { v[ly] = 0.0; }
for (int hy = 0; hy < DHI; ++hy)
{
const real_t zYX = u(lz,hy,hx);
for (int ly = 0; ly < DLO; ++ly) { v[ly] += zYX * Bt(ly,hy); }
}
for (int ly = 0; ly < DLO; ++ly) { u(lz,ly,hx) = v[ly]; }
}
}
MFEM_SYNC_THREAD;
MFEM_FOREACH_THREAD(lz,y,DLO)
{
MFEM_FOREACH_THREAD(ly,x,DLO)
{
for (int lx = 0; lx < DLO; ++lx) { v[lx] = 0.0; }
for (int hx = 0; hx < DHI; ++hx)
{
const real_t zyX = u(lz,ly,hx);
for (int lx = 0; lx < DLO; ++lx) { v[lx] += zyX * Bt(lx,hx); }
}
for (int lx = 0; lx < DLO; ++lx)
{
u_l(lx,ly,lz,e) = v[lx];
}
}
}
});
}
void ProlongationTranspose3D(const int NE, const int D1D, const int Q1D,
const Vector& localH, Vector& localL,
const Array<real_t>& Bt, const Vector& mask)
{
auto x_ = Reshape(localH.Read(), Q1D, Q1D, Q1D, NE);
auto y_ = Reshape(localL.Write(), D1D, D1D, D1D, NE);
@@ -2518,11 +2656,14 @@ void Restriction3D(const int NE, const int D1D, const int Q1D,
}
});
}
} // namespace TransferKernels
void TensorProductPRefinementTransferOperator::Mult(const Vector& x,
Vector& y) const
{
using namespace TransferKernels;
if (lFESpace.GetMesh()->GetNE() == 0)
{
return;
@@ -2531,11 +2672,25 @@ void TensorProductPRefinementTransferOperator::Mult(const Vector& x,
elem_restrict_lex_l->Mult(x, localL);
if (dim == 2)
{
TransferKernels::Prolongation2D(NE, D1D, Q1D, localL, localH, B, mask);
Prolongation2D(NE, D1D, Q1D, localL, localH, B, mask);
}
else if (dim == 3)
{
TransferKernels::Prolongation3D(NE, D1D, Q1D, localL, localH, B, mask);
switch ((D1D << 4 ) | Q1D)
{
case 0x23:
SmemProlongation3D<2,3>(NE, localL, localH, B, mask); break;
case 0x24:
SmemProlongation3D<2,4>(NE, localL, localH, B, mask); break;
case 0x35:
SmemProlongation3D<3,5>(NE, localL, localH, B, mask); break;
case 0x46:
SmemProlongation3D<4,6>(NE, localL, localH, B, mask); break;
case 0x47:
SmemProlongation3D<4,7>(NE, localL, localH, B, mask); break;
default:
Prolongation3D(NE, D1D, Q1D, localL, localH, B, mask); break;
}
}
else
{
@@ -2549,6 +2704,8 @@ void TensorProductPRefinementTransferOperator::Mult(const Vector& x,
void TensorProductPRefinementTransferOperator::MultTranspose(const Vector& x,
Vector& y) const
{
using namespace TransferKernels;
if (lFESpace.GetMesh()->GetNE() == 0)
{
return;
@@ -2557,11 +2714,25 @@ void TensorProductPRefinementTransferOperator::MultTranspose(const Vector& x,
elem_restrict_lex_h->Mult(x, localH);
if (dim == 2)
{
TransferKernels::Restriction2D(NE, D1D, Q1D, localH, localL, Bt, mask);
ProlongationTranspose2D(NE, D1D, Q1D, localH, localL, Bt, mask);
}
else if (dim == 3)
{
TransferKernels::Restriction3D(NE, D1D, Q1D, localH, localL, Bt, mask);
switch ((D1D << 4 ) | Q1D)
{
case 0x23:
SmemProlongationTranspose3D<2,3>(NE, localH, localL, Bt, mask); break;
case 0x24:
SmemProlongationTranspose3D<2,4>(NE, localH, localL, Bt, mask); break;
case 0x35:
SmemProlongationTranspose3D<3,5>(NE, localH, localL, Bt, mask); break;
case 0x46:
SmemProlongationTranspose3D<4,6>(NE, localH, localL, Bt, mask); break;
case 0x47:
SmemProlongationTranspose3D<4,7>(NE, localH, localL, Bt, mask); break;
default:
ProlongationTranspose3D(NE, D1D, Q1D, localH, localL, Bt, mask); break;
}
}
else
{
-3
View File
@@ -621,9 +621,6 @@ public:
const FiniteElementSpace& lFESpace_,
const FiniteElementSpace& hFESpace_);
/// Destructor
virtual ~TensorProductPRefinementTransferOperator() { }
/// @brief Interpolation or prolongation of a vector \p x corresponding to
/// the coarse space to the vector \p y corresponding to the fine space.
void Mult(const Vector& x, Vector& y) const override;
+95 -8
View File
@@ -14,24 +14,111 @@
#include "../config/config.hpp"
#ifdef MFEM_USE_CALIPER
#define MFEM_CONCAT_(X,Y) X##Y
#define MFEM_CONCAT(X,Y) MFEM_CONCAT_(X,Y)
#ifdef MFEM_USE_CALIPER
#include "device.hpp"
#include "backends.hpp"
#ifdef MFEM_USE_MPI
#include "communication.hpp"
#endif
#include <optional>
#include <caliper/cali.h>
#include <caliper/cali-manager.h>
#define MFEM_PERF_FUNCTION CALI_CXX_MARK_FUNCTION
#define MFEM_PERF_BEGIN(s) CALI_MARK_BEGIN(s)
#define MFEM_PERF_END(s) CALI_MARK_END(s)
#endif
namespace mfem
{
namespace internal
{
extern int annotation_sync_stream; // defined in globals.cpp
extern int annotation_sync_mpi; // defined in globals.cpp
extern int annotation_enabled; // defined in globals.cpp
#ifdef MFEM_USE_CALIPER
inline void AnnotationSync()
{
if (!annotation_enabled) { return; }
if (annotation_sync_stream && Device::Allows(Backend::DEVICE_MASK))
{
MFEM_STREAM_SYNC;
}
#ifdef MFEM_USE_MPI
if (annotation_sync_mpi && Mpi::IsInitialized() && !Mpi::IsFinalized())
{
MPI_Barrier(GetGlobalMPI_Comm());
}
#endif
}
struct FunctionAnnotation
{
std::optional<cali::Function> cali_func;
FunctionAnnotation(const char *fname)
{
AnnotationSync();
if (annotation_enabled) { cali_func.emplace(fname); }
}
~FunctionAnnotation() { AnnotationSync(); }
};
struct ScopeAnnotation
{
std::optional<cali::ScopeAnnotation> cali_scope;
ScopeAnnotation(const char *name)
{
AnnotationSync();
if (annotation_enabled) { cali_scope.emplace(name); }
}
~ScopeAnnotation() { AnnotationSync(); }
};
#endif // #ifdef MFEM_USE_CALIPER
} // namespace internal
} // namespace mfem
#ifdef MFEM_USE_CALIPER
#define MFEM_PERF_FUNCTION \
mfem::internal::FunctionAnnotation mfem_func_annotation_(_MFEM_FUNC_NAME)
#define MFEM_PERF_BEGIN(s) \
(mfem::internal::AnnotationSync(), CALI_MARK_BEGIN(s))
#define MFEM_PERF_END(s) \
(mfem::internal::AnnotationSync(), CALI_MARK_END(s))
#define MFEM_PERF_SCOPE(name) \
cali::Annotation::Guard cali_autogenerated_guard_name(cali::Annotation("function").begin(std::string(name).c_str()))
mfem::internal::ScopeAnnotation \
MFEM_CONCAT(mfem_scope_annotation_,__LINE__)(name)
#define MFEM_PERF_SYNC_STREAM(b) (mfem::internal::annotation_sync_stream = (b))
#define MFEM_PERF_SYNC_MPI(b) (mfem::internal::annotation_sync_mpi = (b))
#define MFEM_PERF_SYNC(b) (MFEM_PERF_SYNC_STREAM(b), MFEM_PERF_SYNC_MPI(b))
#define MFEM_PERF_ENABLE (mfem::internal::annotation_enabled = 1)
#define MFEM_PERF_DISABLE (mfem::internal::annotation_enabled = 0)
#else
#else // #ifdef MFEM_USE_CALIPER
#define MFEM_PERF_FUNCTION
#define MFEM_PERF_BEGIN(s)
#define MFEM_PERF_END(s)
#define MFEM_PERF_SCOPE(name)
#endif
#define MFEM_PERF_SYNC_STREAM(b)
#define MFEM_PERF_SYNC_MPI(b)
#define MFEM_PERF_SYNC(b)
#define MFEM_PERF_ENABLE
#define MFEM_PERF_DISABLE
#endif
#endif // #ifdef MFEM_USE_CALIPER
#endif // MFEM_ANNOTATION_HPP
-27
View File
@@ -211,9 +211,6 @@ public:
/// Delete the first entry with value == 'el'.
inline void DeleteFirst(const T &el);
/// Delete entries at @a indices, and resize.
inline void DeleteAt(const Array<int> &indices);
/// Delete the whole array.
inline void DeleteAll();
@@ -938,30 +935,6 @@ inline void Array<T>::DeleteFirst(const T &el)
}
}
template <class T>
inline void Array<T>::DeleteAt(const Array<int> &indices)
{
// Make a copy of the indices, sorted.
Array<int> sorted_indices(indices);
sorted_indices.Sort();
int rm_count = 0;
for (int i = 0; i < size; i++)
{
if (rm_count < sorted_indices.Size() && i == sorted_indices[rm_count])
{
rm_count++;
}
else
{
data[i-rm_count] = data[i]; // shift data rm_count
}
}
// Resize to remove tail
size -= rm_count;
}
template <class T>
inline void Array<T>::DeleteAll()
{
+6 -6
View File
@@ -14,7 +14,7 @@
#include "../config/config.hpp"
#if defined(MFEM_USE_CUDA) && defined(__CUDACC__)
#ifdef MFEM_USE_CUDA
#include <cusparse.h>
#include <library_types.h>
#include <cuda_runtime.h>
@@ -22,7 +22,7 @@
#endif
#include "cuda.hpp"
#if defined(MFEM_USE_HIP) && defined(__HIP__)
#ifdef MFEM_USE_HIP
#include <hip/hip_runtime.h>
#endif
#include "hip.hpp"
@@ -43,7 +43,7 @@
#endif
#endif
#if !defined(MFEM_USE_CUDA_OR_HIP)
#if !(defined(MFEM_USE_CUDA) || defined(MFEM_USE_HIP))
#define MFEM_DEVICE
#define MFEM_HOST
#define MFEM_LAMBDA
@@ -55,7 +55,7 @@
#endif
#if !((defined(MFEM_USE_CUDA) && defined(__CUDA_ARCH__)) || \
(defined(MFEM_USE_HIP) && defined(__HIP_DEVICE_COMPILE__)))
(defined(MFEM_USE_HIP) && defined(__HIP_DEVICE_COMPILE__)))
#define MFEM_SHARED
#define MFEM_SYNC_THREAD
#define MFEM_BLOCK_ID(k) 0
@@ -66,7 +66,7 @@
#endif
// 'double' and 'float' atomicAdd implementation for previous versions of CUDA
#if defined(MFEM_USE_CUDA) && defined(__CUDA_ARCH__) && (__CUDA_ARCH__ < 600)
#if defined(MFEM_USE_CUDA) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 600
MFEM_DEVICE inline mfem::real_t atomicAdd(mfem::real_t *add, mfem::real_t val)
{
unsigned long long int *ptr = (unsigned long long int *) add;
@@ -94,7 +94,7 @@ template <typename T>
MFEM_HOST_DEVICE T AtomicAdd(T &add, const T val)
{
#if ((defined(MFEM_USE_CUDA) && defined(__CUDA_ARCH__)) || \
(defined(MFEM_USE_HIP) && defined(__HIP_DEVICE_COMPILE__)))
(defined(MFEM_USE_HIP) && defined(__HIP_DEVICE_COMPILE__)))
return atomicAdd(&add,val);
#else
T old = add;
+3 -13
View File
@@ -347,23 +347,13 @@ void GroupTopology::Swap(GroupTopology &other)
mfem::Swap(group_mgroup, other.group_mgroup);
}
/// \cond DO_NOT_DOCUMENT
// Initialize the static mpi_type for the specializations of MPITypeMap:
const MPI_Datatype MPITypeMap<bool>::mpi_type = MFEM_MPI_CXX_BOOL;
const MPI_Datatype MPITypeMap<char>::mpi_type = MPI_CHAR;
const MPI_Datatype MPITypeMap<unsigned char>::mpi_type = MPI_UNSIGNED_CHAR;
const MPI_Datatype MPITypeMap<short>::mpi_type = MPI_SHORT;
const MPI_Datatype MPITypeMap<unsigned short>::mpi_type = MPI_UNSIGNED_SHORT;
const MPI_Datatype MPITypeMap<int>::mpi_type = MPI_INT;
const MPI_Datatype MPITypeMap<unsigned int>::mpi_type = MPI_UNSIGNED;
const MPI_Datatype MPITypeMap<long>::mpi_type = MPI_LONG;
const MPI_Datatype MPITypeMap<unsigned long>::mpi_type = MPI_UNSIGNED_LONG;
const MPI_Datatype MPITypeMap<long long>::mpi_type = MPI_LONG_LONG;
const MPI_Datatype MPITypeMap<unsigned long long>::mpi_type =
MPI_UNSIGNED_LONG_LONG;
const MPI_Datatype MPITypeMap<float>::mpi_type = MPI_FLOAT;
const MPI_Datatype MPITypeMap<double>::mpi_type = MPI_DOUBLE;
/// \endcond DO_NOT_DOCUMENT
const MPI_Datatype MPITypeMap<int64_t>::mpi_type = MPI_INT64_T;
const MPI_Datatype MPITypeMap<uint64_t>::mpi_type = MPI_UINT64_T;
GroupCommunicator::GroupCommunicator(const GroupTopology &gt, Mode m)
: gtopo(gt), mode(m)
+20 -63
View File
@@ -23,13 +23,9 @@
#include <mpi.h>
#include <cstdint>
// can't directly use MPI_CXX_BOOL because Microsoft's MPI implementation
// doesn't include MPI_CXX_BOOL. Fallback to MPI_C_BOOL if unavailable.
#ifdef MPI_CXX_BOOL
#define MFEM_MPI_CXX_BOOL MPI_CXX_BOOL
#else
#define MFEM_MPI_CXX_BOOL MPI_C_BOOL
#endif
// Some MPI implementations do not have MPI_CXX_BOOL or do not handle it
// correctly, so we use MPI_UNSIGNED_CHAR as the MPI type for 'bool':
#define MFEM_MPI_CXX_BOOL MPI_UNSIGNED_CHAR
namespace mfem
{
@@ -425,24 +421,16 @@ public:
~GroupCommunicator();
};
/// General MPI message tags used by MFEM
enum MessageTag
{
DEREFINEMENT_MATRIX_CONSTRUCTION_DATA =
291, /// ParFiniteElementSpace ParallelDerefinementMatrix and
/// ParDerefineMatrixOp
};
enum VarMessageTag
{
NEIGHBOR_ELEMENT_RANK_VM, ///< NeighborElementRankMessage
NEIGHBOR_ORDER_VM, ///< NeighborOrderMessage
NEIGHBOR_DEREFINEMENT_VM, ///< NeighborDerefinementMessage
NEIGHBOR_REFINEMENT_VM, ///< NeighborRefinementMessage
NEIGHBOR_PREFINEMENT_VM, ///< NeighborPRefinementMessage
NEIGHBOR_ROW_VM, ///< NeighborRowMessage
REBALANCE_VM, ///< RebalanceMessage
REBALANCE_DOF_VM, ///< RebalanceDofMessage
NEIGHBOR_ELEMENT_RANK_VM, ///< NeighborElementRankMessage
NEIGHBOR_ORDER_VM, ///< NeighborOrderMessage
NEIGHBOR_DEREFINEMENT_VM, ///< NeighborDerefinementMessage
NEIGHBOR_REFINEMENT_VM, ///< NeighborRefinementMessage
NEIGHBOR_PREFINEMENT_VM, ///< NeighborPRefinementMessage
NEIGHBOR_ROW_VM, ///< NeighborRowMessage
REBALANCE_VM, ///< RebalanceMessage
REBALANCE_DOF_VM ///< RebalanceDofMessage
};
/// \brief Variable-length MPI message containing unspecific binary data.
@@ -607,50 +595,10 @@ protected:
template <typename Type> struct MPITypeMap;
// Specializations of MPITypeMap; mpi_type initialized in communication.cpp:
template<> struct MPITypeMap<bool>
{
static MFEM_EXPORT const MPI_Datatype mpi_type;
};
template<> struct MPITypeMap<char>
{
static MFEM_EXPORT const MPI_Datatype mpi_type;
};
template<> struct MPITypeMap<unsigned char>
{
static MFEM_EXPORT const MPI_Datatype mpi_type;
};
template<> struct MPITypeMap<short>
{
static MFEM_EXPORT const MPI_Datatype mpi_type;
};
template<> struct MPITypeMap<unsigned short>
{
static MFEM_EXPORT const MPI_Datatype mpi_type;
};
template<> struct MPITypeMap<int>
{
static MFEM_EXPORT const MPI_Datatype mpi_type;
};
template<> struct MPITypeMap<unsigned int>
{
static MFEM_EXPORT const MPI_Datatype mpi_type;
};
template<> struct MPITypeMap<long>
{
static MFEM_EXPORT const MPI_Datatype mpi_type;
};
template<> struct MPITypeMap<unsigned long>
{
static MFEM_EXPORT const MPI_Datatype mpi_type;
};
template<> struct MPITypeMap<long long>
{
static MFEM_EXPORT const MPI_Datatype mpi_type;
};
template<> struct MPITypeMap<unsigned long long>
{
static MFEM_EXPORT const MPI_Datatype mpi_type;
};
template<> struct MPITypeMap<double>
{
static MFEM_EXPORT const MPI_Datatype mpi_type;
@@ -659,6 +607,15 @@ template<> struct MPITypeMap<float>
{
static MFEM_EXPORT const MPI_Datatype mpi_type;
};
template<> struct MPITypeMap<int64_t>
{
static MFEM_EXPORT const MPI_Datatype mpi_type;
};
template<> struct MPITypeMap<uint64_t>
{
static MFEM_EXPORT const MPI_Datatype mpi_type;
};
/** Reorder MPI ranks to follow the Z-curve within the physical machine topology
(provided that functions to query physical node coordinates are available).
+1 -1
View File
@@ -24,7 +24,7 @@ void mfem_cuda_error(cudaError_t err, const char *expr, const char *func,
const char *file, int line)
{
mfem::err << "\n\nCUDA error: (" << expr << ") failed with error:\n --> "
<< cudaGetErrorString(err) << " [code: " << (int)err << ']'
<< cudaGetErrorString(err)
<< "\n ... in function: " << func
<< "\n ... in file: " << file << ':' << line << '\n';
mfem_error();
+5 -5
View File
@@ -18,7 +18,7 @@
// CUDA block size used by MFEM.
#define MFEM_CUDA_BLOCKS 256
#if defined(MFEM_USE_CUDA) && defined(__CUDACC__)
#ifdef MFEM_USE_CUDA
#define MFEM_USE_CUDA_OR_HIP
#define MFEM_DEVICE __device__
#define MFEM_HOST __host__
@@ -37,9 +37,10 @@
__FILE__, __LINE__); \
} \
} while (0)
#endif // MFEM_USE_CUDA
// Define the MFEM inner threading macros
#if defined(__CUDA_ARCH__)
#if defined(MFEM_USE_CUDA) && defined(__CUDA_ARCH__)
#define MFEM_SHARED __shared__
#define MFEM_SYNC_THREAD __syncthreads()
#define MFEM_BLOCK_ID(k) blockIdx.k
@@ -47,13 +48,12 @@
#define MFEM_THREAD_SIZE(k) blockDim.k
#define MFEM_FOREACH_THREAD(i,k,N) for(int i=threadIdx.k; i<N; i+=blockDim.k)
#define MFEM_FOREACH_THREAD_DIRECT(i,k,N) if(const int i=threadIdx.k; i<N)
#endif // defined(__CUDA_ARCH__)
#endif // defined(MFEM_USE_CUDA) && defined(__CUDACC__)
#endif
namespace mfem
{
#if defined(MFEM_USE_CUDA) && defined(__CUDACC__)
#ifdef MFEM_USE_CUDA
// Function used by the macro MFEM_GPU_CHECK.
void mfem_cuda_error(cudaError_t err, const char *expr, const char *func,
const char *file, int line);
+17 -1
View File
@@ -151,6 +151,22 @@ Device::Device()
{
SetGPUAwareMPI(true);
}
if (const char *mfem_perf_sync = GetEnv("MFEM_PERF_SYNC"))
{
MFEM_PERF_SYNC(std::atoi(mfem_perf_sync));
MFEM_CONTRACT_VAR(mfem_perf_sync);
}
if (const char *mfem_perf_sync_stream = GetEnv("MFEM_PERF_SYNC_STREAM"))
{
MFEM_PERF_SYNC_STREAM(std::atoi(mfem_perf_sync_stream));
MFEM_CONTRACT_VAR(mfem_perf_sync_stream);
}
if (const char *mfem_perf_sync_mpi = GetEnv("MFEM_PERF_SYNC_MPI"))
{
MFEM_PERF_SYNC_MPI(std::atoi(mfem_perf_sync_mpi));
MFEM_CONTRACT_VAR(mfem_perf_sync_mpi);
}
}
Device::~Device()
@@ -615,7 +631,7 @@ void Device::Setup(const std::string &device_option, const int device_id)
if (Allows(Backend::DEBUG_DEVICE)) { ngpu = 1; }
}
MemoryType Device::QueryMemoryType(const void* ptr)
MemoryType Device::QueryMemoryType(void *ptr)
{
// from HYPRE's hypre_GetPointerLocation
MemoryType res = MemoryType::HOST;
+3 -3
View File
@@ -297,9 +297,9 @@ public:
/// Get the status of GPU-aware MPI flag.
static bool GetGPUAwareMPI() { return Get().mpi_gpu_aware; }
/** Query the device driver for what memory type a given @a ptr is allocated
* with. */
static MemoryType QueryMemoryType(const void* ptr);
/** @brief Query the device driver for what memory type a given @a ptr is
allocated with. */
static MemoryType QueryMemoryType(void *ptr);
/** @brief The number of hardware compute units/streaming multiprocessors
available on a given compute device @a device_id. */
+1 -1
View File
@@ -193,4 +193,4 @@ void mfem_warning(const char *msg)
}
}
}
} // namespace mfem
+2 -2
View File
@@ -176,7 +176,7 @@ __device__ void abort_msg(T & msg)
printf(__VA_ARGS__); \
asm("trap;"); \
}
#elif defined(__HIP_DEVICE_COMPILE__)
#elif defined(MFEM_USE_HIP)
#define MFEM_ABORT_KERNEL(...) \
{ \
printf(__VA_ARGS__); \
@@ -208,4 +208,4 @@ __device__ void abort_msg(T & msg)
#define MFEM_ASSERT_KERNEL(x,...)
#endif
#endif
#endif // MFEM_ERROR_HPP
+12 -171
View File
@@ -158,8 +158,8 @@ private:
#define MFEM_PRAGMA(X) _Pragma(#X)
// MFEM_UNROLL pragma macro that can be used inside MFEM_FORALL macros.
#if defined(MFEM_USE_CUDA) && defined(__CUDA_ARCH__) // Clang cuda or nvcc
#ifdef __NVCC__ // nvcc specifically
#if defined(MFEM_USE_CUDA) && defined(__CUDA_ARCH__)
#ifdef __NVCC__
#define MFEM_UNROLL(N) MFEM_PRAGMA(unroll(N))
#else // Assuming Clang CUDA
#define MFEM_UNROLL(N) MFEM_PRAGMA(unroll N)
@@ -169,12 +169,12 @@ private:
#endif
// MFEM_GPU_FORALL: "parallel for" executed with CUDA or HIP based on the MFEM
// build-time configuration (MFEM_USE_CUDA or MFEM_USE_HIP), and if compiling
// with CUDA/HIP language. Otherwise, this macro is a no-op.
#if defined(MFEM_USE_CUDA) && defined(__CUDACC__)
// build-time configuration (MFEM_USE_CUDA or MFEM_USE_HIP). If neither CUDA nor
// HIP is enabled, this macro is a no-op.
#if defined(MFEM_USE_CUDA)
#define MFEM_GPU_FORALL(i, N,...) CuWrap1D(N, [=] MFEM_DEVICE \
(int i) {__VA_ARGS__})
#elif defined(MFEM_USE_HIP) && defined(__HIP__)
#elif defined(MFEM_USE_HIP)
#define MFEM_GPU_FORALL(i, N,...) HipWrap1D(N, [=] MFEM_DEVICE \
(int i) {__VA_ARGS__})
#else
@@ -225,52 +225,6 @@ void OmpWrap(const int N, HBODY &&h_body)
#endif
}
template <typename HBODY>
void OmpWrap2D(const int Nx, const int Ny, HBODY &&h_body)
{
#ifdef MFEM_USE_OPENMP
// requires OpenMP 3.1
#pragma omp parallel for collapse(2)
for (int j = 0; j < Ny; j++)
{
for (int i = 0; i < Nx; i++)
{
h_body(i, j);
}
}
#else
MFEM_CONTRACT_VAR(Nx);
MFEM_CONTRACT_VAR(Ny);
MFEM_CONTRACT_VAR(h_body);
MFEM_ABORT("OpenMP requested for MFEM but OpenMP is not enabled!");
#endif
}
template <typename HBODY>
void OmpWrap3D(const int Nx, const int Ny, const int Nz, HBODY &&h_body)
{
#ifdef MFEM_USE_OPENMP
// requires OpenMP 3.1
#pragma omp parallel for collapse(3)
for (int k = 0; k < Nz; k++)
{
for (int j = 0; j < Ny; j++)
{
for (int i = 0; i < Nx; i++)
{
h_body(i, j, k);
}
}
}
#else
MFEM_CONTRACT_VAR(Nx);
MFEM_CONTRACT_VAR(Ny);
MFEM_CONTRACT_VAR(Nz);
MFEM_CONTRACT_VAR(h_body);
MFEM_ABORT("OpenMP requested for MFEM but OpenMP is not enabled!");
#endif
}
/// RAJA Cuda and Hip backends
#if defined(MFEM_USE_RAJA) && defined(RAJA_ENABLE_CUDA)
@@ -500,42 +454,6 @@ void RajaOmpWrap(const int N, HBODY &&h_body)
RAJA::forall<RAJA::omp_parallel_for_exec>(RAJA::RangeSegment(0,N), h_body);
}
template <typename HBODY>
void RajaOmpWrap2D(const int Nx, const int Ny, HBODY &&h_body)
{
using omp_launch_policy = RAJA::LaunchPolicy<RAJA::omp_launch_t>;
using global_thread_xy = RAJA::LoopPolicy<RAJA::omp_for_exec>;
RAJA::RangeSegment xrange(0, Nx);
RAJA::RangeSegment yrange(0, Ny);
RAJA::launch<omp_launch_policy>(RAJA::ExecPlace::HOST, RAJA::LaunchParams(),
[=](RAJA::LaunchContext ctx)
{
// contiguous in x
RAJA::expt::loop<global_thread_xy>(ctx, xrange, yrange, [&](int i, int j)
{
h_body(i, j);
});
});
}
template <typename HBODY>
void RajaOmpWrap3D(const int Nx, const int Ny, const int Nz, HBODY &&h_body)
{
using omp_launch_policy = RAJA::LaunchPolicy<RAJA::omp_launch_t>;
using global_thread_xyz = RAJA::LoopPolicy<RAJA::omp_for_exec>;
RAJA::RangeSegment xrange(0, Nx);
RAJA::RangeSegment yrange(0, Ny);
RAJA::RangeSegment zrange(0, Nz);
RAJA::launch<omp_launch_policy>(RAJA::ExecPlace::HOST, RAJA::LaunchParams(),
[=](RAJA::LaunchContext ctx)
{
// contiguous in x
RAJA::expt::loop<global_thread_xyz>(ctx, xrange, yrange, zrange,
[&](int i, int j, int k)
{ h_body(i, j, k); });
});
}
#endif
@@ -563,7 +481,7 @@ void RajaSeqWrap(const int N, HBODY &&h_body)
/// CUDA backend
#if defined(MFEM_USE_CUDA) && defined(__CUDACC__)
#ifdef MFEM_USE_CUDA
template <typename BODY> __global__ static
void CuKernel1D(const int N, BODY body)
@@ -655,11 +573,11 @@ struct CuWrap<3>
}
};
#endif // defined(MFEM_USE_CUDA) && defined(__CUDACC__)
#endif // MFEM_USE_CUDA
/// HIP backend
#if defined(MFEM_USE_HIP) && defined(__HIP__)
#ifdef MFEM_USE_HIP
template <typename BODY> __global__ static
void HipKernel1D(const int N, BODY body)
@@ -750,7 +668,7 @@ struct HipWrap<3>
}
};
#endif // defined(MFEM_USE_HIP) && defined(__HIP__)
#endif // MFEM_USE_HIP
/// The forall kernel body wrapper
@@ -783,7 +701,7 @@ inline void ForallWrap(const bool use_dev, const int N,
}
#endif
#if defined(MFEM_USE_CUDA) && defined(__CUDACC__)
#ifdef MFEM_USE_CUDA
// If Backend::CUDA is allowed, use it
if (Device::Allows(Backend::CUDA))
{
@@ -791,7 +709,7 @@ inline void ForallWrap(const bool use_dev, const int N,
}
#endif
#if defined(MFEM_USE_HIP) && defined(__HIP__)
#ifdef MFEM_USE_HIP
// If Backend::HIP is allowed, use it
if (Device::Allows(Backend::HIP))
{
@@ -835,83 +753,6 @@ inline void ForallWrap(const bool use_dev, const int N, lambda &&body,
template<typename lambda>
inline void forall(int N, lambda &&body) { ForallWrap<1>(true, N, body); }
template<typename lambda>
inline void forall(int Nx, int Ny, lambda &&body)
{
if (Device::Allows(Backend::DEVICE_MASK))
{
forall(Nx * Ny, [=] MFEM_HOST_DEVICE(int idx)
{
int j = idx / Nx;
int i = idx % Nx;
body(i, j);
});
}
#if defined(MFEM_USE_RAJA) && defined(RAJA_ENABLE_OPENMP)
else if (Device::Allows(Backend::RAJA_OMP))
{
return RajaOmpWrap2D(Nx, Ny, body);
}
#endif
#ifdef MFEM_USE_OPENMP
else if (Device::Allows(Backend::OMP))
{
return OmpWrap2D(Nx, Ny, body);
}
#endif
else
{
for (int j = 0; j < Ny; ++j)
{
for (int i = 0; i < Nx; ++i)
{
body(i, j);
}
}
}
}
template<typename lambda>
inline void forall(int Nx, int Ny, int Nz, lambda &&body)
{
if (Device::Allows(Backend::DEVICE_MASK))
{
forall(Nx * Ny * Nz, [=] MFEM_HOST_DEVICE(int idx)
{
int i = idx % Nx;
int j = idx / Nx;
int k = j / Ny;
j = j % Ny;
body(i, j, k);
});
}
#if defined(MFEM_USE_RAJA) && defined(RAJA_ENABLE_OPENMP)
else if (Device::Allows(Backend::RAJA_OMP))
{
return RajaOmpWrap3D(Nx, Ny, Nz, body);
}
#endif
#ifdef MFEM_USE_OPENMP
else if (Device::Allows(Backend::OMP))
{
return OmpWrap3D(Nx, Ny, Nz, body);
}
#endif
else
{
for (int k = 0; k < Nz; ++k)
{
for (int j = 0; j < Ny; ++j)
{
for (int i = 0; i < Nx; ++i)
{
body(i, j, k);
}
}
}
}
}
template<typename lambda>
inline void forall_switch(bool use_dev, int N, lambda &&body)
{
+4
View File
@@ -31,6 +31,10 @@ namespace internal
{
bool mfem_out_initialized = false;
bool mfem_err_initialized = false;
int annotation_sync_stream = 0; // declared in annotation.hpp
int annotation_sync_mpi = 0; // declared in annotation.hpp
int annotation_enabled = 1; // declared in annotation.hpp
}
void OutStream::Init()
+1 -1
View File
@@ -24,7 +24,7 @@ void mfem_hip_error(hipError_t err, const char *expr, const char *func,
const char *file, int line)
{
mfem::err << "\n\nHIP error: (" << expr << ") failed with error:\n --> "
<< hipGetErrorString(err) << " [code: " << (int)err << ']'
<< hipGetErrorString(err)
<< "\n ... in function: " << func
<< "\n ... in file: " << file << ':' << line << '\n';
mfem_error();
+4 -4
View File
@@ -18,7 +18,7 @@
// HIP block size used by MFEM.
#define MFEM_HIP_BLOCKS 256
#if defined(MFEM_USE_HIP) && defined(__HIP__)
#ifdef MFEM_USE_HIP
#define MFEM_USE_CUDA_OR_HIP
#define MFEM_DEVICE __device__
#define MFEM_HOST __host__
@@ -37,9 +37,10 @@
__FILE__, __LINE__); \
} \
} while (0)
#endif // MFEM_USE_HIP
// Define the MFEM inner threading macros
#if defined(__HIP_DEVICE_COMPILE__)
#if defined(MFEM_USE_HIP) && defined(__HIP_DEVICE_COMPILE__)
#define MFEM_SHARED __shared__
#define MFEM_SYNC_THREAD __syncthreads()
#define MFEM_BLOCK_ID(k) hipBlockIdx_ ##k
@@ -49,8 +50,7 @@
for(int i=hipThreadIdx_ ##k; i<N; i+=hipBlockDim_ ##k)
#define MFEM_FOREACH_THREAD_DIRECT(i,k,N) \
if(const int i=hipThreadIdx_ ##k; i<N)
#endif // defined(__HIP_DEVICE_COMPILE__)
#endif // defined(MFEM_USE_HIP) && defined(__HIP__)
#endif
namespace mfem
{
+8 -2
View File
@@ -513,7 +513,10 @@ public:
void *HtoD(void *dst, const void *src, size_t bytes) override
{ return HipMemcpyHtoD(dst, src, bytes); }
void *DtoD(void* dst, const void* src, size_t bytes) override
{ return HipMemcpyDtoD(dst, src, bytes); }
// Unlike cudaMemcpy(DtoD), hipMemcpy(DtoD) causes a host-side synchronization so
// instead we use hipMemcpyAsync to get similar behavior.
// for more info see: https://github.com/mfem/mfem/pull/2780
{ return HipMemcpyDtoDAsync(dst, src, bytes); }
void *DtoH(void *dst, const void *src, size_t bytes) override
{ return HipMemcpyDtoH(dst, src, bytes); }
};
@@ -655,7 +658,10 @@ public:
return CuMemcpyDtoD(dst, src, bytes);
#endif
#ifdef MFEM_USE_HIP
return HipMemcpyDtoD(dst, src, bytes);
// Unlike cudaMemcpy(DtoD), hipMemcpy(DtoD) causes a host-side synchronization so
// instead we use hipMemcpyAsync to get similar behavior.
// for more info see: https://github.com/mfem/mfem/pull/2780
return HipMemcpyDtoDAsync(dst, src, bytes);
#endif
// rm.copy(dst, const_cast<void*>(src), bytes); return dst;
}
+7 -7
View File
@@ -657,7 +657,8 @@ private: // Static methods used by the Memory<T> class
/// Return the host pointer.
MFEM_ENZYME_INACTIVE static void *Register_(void *ptr, void *h_ptr,
size_t bytes, MemoryType mt,
bool own, bool alias, unsigned &flags);
bool own, bool alias,
unsigned &flags);
/// Register a pair of external host and device pointers
static void Register2_(void *h_ptr, void *d_ptr, size_t bytes,
@@ -741,7 +742,7 @@ private:
/// Insert a host address @a h_ptr and size *a bytes in the memory map to be
/// managed.
void Insert(void *h_ptr, size_t bytes, MemoryType h_mt, MemoryType d_mt);
void Insert(void *h_ptr, size_t bytes, MemoryType h_mt, MemoryType d_mt);
/// Insert a device and the host addresses in the memory map
void InsertDevice(void *d_ptr, void *h_ptr, size_t bytes,
@@ -896,7 +897,6 @@ inline HYPRE_MemoryLocation GetHypreMemoryLocation()
#elif MFEM_HYPRE_VERSION < 23100
return HYPRE_MEMORY_DEVICE;
#else // HYPRE_USING_GPU is defined and MFEM_HYPRE_VERSION >= 23100
if (!HYPRE_Initialized()) { return HYPRE_MEMORY_HOST; }
HYPRE_MemoryLocation loc;
HYPRE_GetMemoryLocation(&loc);
return loc;
@@ -981,7 +981,7 @@ inline void Memory<T>::Wrap(T *ptr, int size, bool own)
#ifdef MFEM_DEBUG
if (own && MemoryManager::Exists())
{
MemoryType h_ptr_mt = MemoryManager::GetHostMemoryType_(h_ptr);
MemoryType h_ptr_mt = MemoryManager::GetHostMemoryType_((void*)h_ptr);
MFEM_VERIFY(h_mt == h_ptr_mt,
"h_mt = " << (int)h_mt << ", h_ptr_mt = " << (int)h_ptr_mt);
}
@@ -989,7 +989,8 @@ inline void Memory<T>::Wrap(T *ptr, int size, bool own)
if (own && h_mt != MemoryType::HOST)
{
const size_t bytes = size*sizeof(T);
MemoryManager::Register_(ptr, ptr, bytes, h_mt, own, false, flags);
MemoryManager::Register_((void*)ptr, (void*)ptr, bytes, h_mt, own, false,
flags);
}
}
@@ -1058,8 +1059,7 @@ inline void Memory<T>::MakeAlias(const Memory &base, int offset, int size)
// register the 'base' if the MemoryManager::Exists():
MemoryManager::Exists()
#else // HYPRE_USING_GPU is defined and MFEM_HYPRE_VERSION >= 23100
IsDeviceMemory(MemoryManager::GetDeviceMemoryType()) ||
(MemoryManager::Exists() && HypreUsingGPU())
MemoryManager::Exists() && HypreUsingGPU()
#endif
)
{
+1 -1
View File
@@ -537,7 +537,7 @@ void reduce(int N, T &res, B &&body, const R &reducer, bool use_dev,
return;
}
#if defined(MFEM_USE_CUDA_OR_HIP)
#if defined(MFEM_USE_HIP) || defined(MFEM_USE_CUDA)
if (use_dev &&
mfem::Device::Allows(Backend::CUDA | Backend::HIP | Backend::RAJA_CUDA |
Backend::RAJA_HIP))
-176
View File
@@ -1,176 +0,0 @@
// Copyright (c) 2010-2025, Lawrence Livermore National Security, LLC. Produced
// at the Lawrence Livermore National Laboratory. All Rights reserved. See files
// LICENSE and NOTICE for details. LLNL-CODE-806117.
//
// This file is part of the MFEM library. For more information and source code
// availability visit https://mfem.org.
//
// MFEM is free software; you can redistribute it and/or modify it under the
// terms of the BSD-3 license. We welcome feedback and contributions, see file
// CONTRIBUTING.md for details.
#ifndef MFEM_SCAN_HPP
#define MFEM_SCAN_HPP
#ifdef MFEM_USE_CUDA
#include <cub/device/device_scan.cuh>
#define MFEM_CUB_NAMESPACE cub
#elif MFEM_USE_HIP
#include <hipcub/device/device_scan.hpp>
#define MFEM_CUB_NAMESPACE hipcub
#endif
#include <functional>
#include <numeric>
namespace mfem
{
/// Equivalent to InclusiveScan(use_dev, d_in, d_out, num_items, workspace,
/// std::plus<>{})
template <class InputIt, class OutputIt>
void InclusiveScan(bool use_dev, InputIt d_in, OutputIt d_out, size_t num_items,
Array<char> &workspace)
{
// forward to InclusiveSum for potentially faster kernels
#if defined(MFEM_USE_CUDA) || defined(MFEM_USE_HIP)
if (use_dev && mfem::Device::Allows(Backend::CUDA_MASK | Backend::HIP_MASK))
{
size_t bytes = workspace.Size();
if (bytes)
{
auto err = MFEM_CUB_NAMESPACE::DeviceScan::InclusiveSum(
workspace.Write(), bytes, d_in, d_out, num_items);
#if defined(MFEM_USE_CUDA)
if (err == cudaSuccess)
{
return;
}
#elif defined(MFEM_USE_HIP)
if (err == hipSuccess)
{
return;
}
#endif
}
// try allocating a larger buffer
bytes = 0;
MFEM_GPU_CHECK(MFEM_CUB_NAMESPACE::DeviceScan::InclusiveSum(
nullptr, bytes, d_in, d_out, num_items));
workspace.SetSize(bytes);
MFEM_GPU_CHECK(MFEM_CUB_NAMESPACE::DeviceScan::InclusiveSum(
workspace.Write(), bytes, d_in, d_out, num_items));
return;
}
#endif
std::inclusive_scan(d_in, d_in + num_items, d_out);
}
/// Performs an inclusive scan of [d_in, d_in+num_items) -> [d_out,
/// d_out+num_items). This call is potentially asynchronous on the device.
/// @a d_in input start.
/// @a d_out output start. Can perform in-place scans with d_out = d_in
/// @a workspace temporary workspace used for device scans. TODO: replace with
/// internal temporary workspace once that's added to the memory manager.
/// @a scan_op binary scan functor. Must be associative. If only weakly
/// associative (i.e. floating point addition) results are not deterministic. On
/// device this must also be commutative.
template <class InputIt, class OutputIt, class ScanOp>
void InclusiveScan(bool use_dev, InputIt d_in, OutputIt d_out, size_t num_items,
Array<char> &workspace, ScanOp scan_op)
{
#if defined(MFEM_USE_CUDA) || defined(MFEM_USE_HIP)
if (use_dev && mfem::Device::Allows(Backend::CUDA_MASK | Backend::HIP_MASK))
{
size_t bytes = workspace.Size();
if (bytes)
{
auto err = MFEM_CUB_NAMESPACE::DeviceScan::InclusiveScan(
workspace.Write(), bytes, d_in, d_out, scan_op, num_items);
#if defined(MFEM_USE_CUDA)
if (err == cudaSuccess)
{
return;
}
#elif defined(MFEM_USE_HIP)
if (err == hipSuccess)
{
return;
}
#endif
}
// try allocating a larger buffer
bytes = 0;
MFEM_GPU_CHECK(MFEM_CUB_NAMESPACE::DeviceScan::InclusiveScan(
nullptr, bytes, d_in, d_out, scan_op, num_items));
workspace.SetSize(bytes);
MFEM_GPU_CHECK(MFEM_CUB_NAMESPACE::DeviceScan::InclusiveScan(
workspace.Write(), bytes, d_in, d_out, scan_op, num_items));
return;
}
#endif
std::inclusive_scan(d_in, d_in + num_items, d_out, scan_op);
}
/// Performs an exclusive scan of [d_in, d_in+num_items) -> [d_out,
/// d_out+num_items). This call is potentially asynchronous on the device.
/// @a d_in input start.
/// @a d_out output start. Can perform in-place scans with d_out = d_in
/// @a workspace temporary workspace used for device scans. TODO: replace with
/// internal temporary workspace once that's added to the memory manager.
/// @a scan_op binary scan functor. Must be associative. If only weakly
/// associative (i.e. floating point addition) results are not deterministic. On
/// device this must also be commutative.
template <class InputIt, class OutputIt, class T, class ScanOp>
void ExclusiveScan(bool use_dev, InputIt d_in, OutputIt d_out, size_t num_items,
T init_value, Array<char> &workspace, ScanOp scan_op)
{
#if defined(MFEM_USE_CUDA) || defined(MFEM_USE_HIP)
if (use_dev && mfem::Device::Allows(Backend::CUDA_MASK | Backend::HIP_MASK))
{
size_t bytes = workspace.Size();
if (bytes)
{
auto err = MFEM_CUB_NAMESPACE::DeviceScan::ExclusiveScan(
workspace.Write(), bytes, d_in, d_out, scan_op, init_value,
num_items);
#if defined(MFEM_USE_CUDA)
if (err == cudaSuccess)
{
return;
}
#elif defined(MFEM_USE_HIP)
if (err == hipSuccess)
{
return;
}
#endif
}
// try allocating a larger buffer
bytes = 0;
MFEM_GPU_CHECK(MFEM_CUB_NAMESPACE::DeviceScan::ExclusiveScan(
nullptr, bytes, d_in, d_out, scan_op, init_value, num_items));
workspace.SetSize(bytes);
MFEM_GPU_CHECK(MFEM_CUB_NAMESPACE::DeviceScan::ExclusiveScan(
workspace.Write(), bytes, d_in, d_out, scan_op, init_value,
num_items));
return;
}
#endif
std::exclusive_scan(d_in, d_in + num_items, d_out, init_value, scan_op);
}
/// Equivalent to ExclusiveScan(use_dev, d_in, d_out, num_items, init_value,
/// workspace, std::plus<>{})
template <class InputIt, class OutputIt, class T>
void ExclusiveScan(bool use_dev, InputIt d_in, OutputIt d_out, size_t num_items,
T init_value, Array<char> &workspace)
{
ExclusiveScan(use_dev, d_in, d_out, num_items, init_value, workspace,
std::plus<> {});
}
} // namespace mfem
#undef MFEM_CUB_NAMESPACE
#endif
+4 -4
View File
@@ -20,13 +20,13 @@
#define MFEM_CU_or_HIP(stub) HIP##stub
#endif
#define MFEM_CONCAT(x, y, z) MFEM_CONCAT_(x, y, z)
#define MFEM_CONCAT_(x, y, z) x ## y ## z
#define MFEM_CONCAT3(x, y, z) MFEM_CONCAT3_(x, y, z)
#define MFEM_CONCAT3_(x, y, z) x ## y ## z
#ifdef MFEM_USE_SINGLE
#define MFEM_GPUBLAS_PREFIX(stub) MFEM_CONCAT(MFEM_cu_or_hip(blas), S, stub)
#define MFEM_GPUBLAS_PREFIX(stub) MFEM_CONCAT3(MFEM_cu_or_hip(blas), S, stub)
#elif defined(MFEM_USE_DOUBLE)
#define MFEM_GPUBLAS_PREFIX(stub) MFEM_CONCAT(MFEM_cu_or_hip(blas), D, stub)
#define MFEM_GPUBLAS_PREFIX(stub) MFEM_CONCAT3(MFEM_cu_or_hip(blas), D, stub)
#endif
#define MFEM_BLAS_SUCCESS MFEM_CU_or_HIP(BLAS_STATUS_SUCCESS)
-28
View File
@@ -4405,32 +4405,4 @@ void BatchLUSolve(const DenseTensor &Mlu, const Array<int> &P, Vector &X)
BatchedLinAlg::LUSolve(Mlu, P, X);
}
#ifdef MFEM_USE_LAPACK
void BandedSolve(int KL, int KU, DenseMatrix &AB, DenseMatrix &B,
Array<int> &ipiv)
{
int LDAB = (2*KL) + KU + 1;
int N = AB.NumCols();
int NRHS = B.NumCols();
int info;
ipiv.SetSize(N);
MFEM_LAPACK_PREFIX(gbsv_)(&N, &KL, &KU, &NRHS, AB.GetData(), &LDAB,
ipiv.GetData(), B.GetData(), &N, &info);
MFEM_ASSERT(info == 0, "BandedSolve failed in LAPACK");
}
void BandedFactorizedSolve(int KL, int KU, DenseMatrix &AB, DenseMatrix &B,
bool transpose, Array<int> &ipiv)
{
int LDAB = (2*KL) + KU + 1;
int N = AB.NumCols();
int NRHS = B.NumCols();
char trans = transpose ? 'T' : 'N';
int info;
MFEM_LAPACK_PREFIX(gbtrs_)(&trans, &N, &KL, &KU, &NRHS, AB.GetData(), &LDAB,
ipiv.GetData(), B.GetData(), &N, &info);
MFEM_ASSERT(info == 0, "BandedFactorizedSolve failed in LAPACK");
}
#endif
} // namespace mfem
-7
View File
@@ -1329,13 +1329,6 @@ void BatchLUFactor(DenseTensor &Mlu, Array<int> &P, const real_t TOL = 0.0);
dimension m x n. */
void BatchLUSolve(const DenseTensor &Mlu, const Array<int> &P, Vector &X);
#ifdef MFEM_USE_LAPACK
void BandedSolve(int KL, int KU, DenseMatrix &AB, DenseMatrix &B,
Array<int> &ipiv);
void BandedFactorizedSolve(int KL, int KU, DenseMatrix &AB, DenseMatrix &B,
bool transpose, Array<int> &ipiv);
#endif
// Inline methods
inline real_t &DenseMatrix::operator()(int i, int j)
+4
View File
@@ -4091,6 +4091,8 @@ void HypreSolver::Setup(const HypreParVector &b, HypreParVector &x) const
{
if (setup_called) { return; }
MFEM_PERF_FUNCTION;
MFEM_VERIFY(A != NULL, "HypreParMatrix A is missing");
HYPRE_Int err_flag = SetupFcn()(*this, *A, b, x);
@@ -4116,6 +4118,8 @@ void HypreSolver::Setup(const Vector &b, Vector &x) const
void HypreSolver::Mult(const HypreParVector &b, HypreParVector &x) const
{
MFEM_PERF_FUNCTION;
HYPRE_Int err_flag;
if (A == NULL)
{
-7
View File
@@ -42,13 +42,6 @@ extern "C" void
MFEM_LAPACK_PREFIX(getri_)(int *N, real_t *A, int *LDA, int *IPIV, real_t *WORK,
int *LWORK, int *INFO);
extern "C" void
MFEM_LAPACK_PREFIX(gbsv_)(int *, int *, int *, int *, real_t *, int *, int *,
real_t *, int *, int *);
extern "C" void
MFEM_LAPACK_PREFIX(gbtrs_)(char *, int *, int *, int *, int *, real_t *, int *,
int *, real_t *, int *, int *);
extern "C" void
MFEM_LAPACK_PREFIX(syevr_)(char *JOBZ, char *RANGE, char *UPLO, int *N,
real_t *A, int *LDA, real_t *VL, real_t *VU, int *IL,
int *IU, real_t *ABSTOL, int *M, real_t *W,
+8 -6
View File
@@ -50,17 +50,19 @@ void Operator::InitTVectors(const Operator *Po, const Operator *Ri,
void Operator::AddMult(const Vector &x, Vector &y, const real_t a) const
{
mfem::Vector z(y.Size());
Mult(x, z);
y.Add(a, z);
z_am.SetSize(y.Size());
z_am.UseDevice(true);
Mult(x, z_am);
y.Add(a, z_am);
}
void Operator::AddMultTranspose(const Vector &x, Vector &y,
const real_t a) const
{
mfem::Vector z(y.Size());
MultTranspose(x, z);
y.Add(a, z);
z_am.SetSize(y.Size());
z_am.UseDevice(true);
MultTranspose(x, z_am);
y.Add(a, z_am);
}
void Operator::ArrayMult(const Array<const Vector *> &X,
+7
View File
@@ -23,6 +23,13 @@ class RectangularConstrainedOperator;
/// Abstract operator
class Operator
{
private:
/// Auxiliary Vector used by the methods AddMult() and AddMultTranspose().
/** @note This Vector is private to prevent derived classes from accidentaly
using it in their implementation of Mult() or MultTranspose() which may
lead to hard-to-find bugs. */
mutable Vector z_am;
protected:
int height; ///< Dimension of the output / number of rows in the matrix.
int width; ///< Dimension of the input / number of columns in the matrix.
+79 -36
View File
@@ -314,25 +314,29 @@ void OperatorJacobiSmoother::Mult(const Vector &x, Vector &y) const
MFEM_VERIFY(x.Size() == Width(), "invalid input vector");
MFEM_VERIFY(y.Size() == Height(), "invalid output vector");
auto DI = dinv.Read();
auto X = x.Read();
if (iterative_mode)
{
MFEM_VERIFY(oper, "iterative_mode == true requires the forward operator");
oper->Mult(y, residual); // r = A y
subtract(x, residual, residual); // r = x - A y
auto R = residual.Read();
auto Y = y.ReadWrite();
// y += D^{-1} (x - A y)
mfem::forall(height, [=] MFEM_HOST_DEVICE (int i)
{
Y[i] += DI[i] * (X[i] - R[i]);
});
}
else
{
residual = x;
y.UseDevice(true);
y = 0.0;
auto Y = y.Write();
// y = D^{-1} x
mfem::forall(height, [=] MFEM_HOST_DEVICE (int i)
{
Y[i] = DI[i] * X[i];
});
}
auto DI = dinv.Read();
auto R = residual.Read();
auto Y = y.ReadWrite();
mfem::forall(height, [=] MFEM_HOST_DEVICE (int i)
{
Y[i] += DI[i] * R[i];
});
}
OperatorChebyshevSmoother::OperatorChebyshevSmoother(const Operator &oper_,
@@ -348,7 +352,8 @@ OperatorChebyshevSmoother::OperatorChebyshevSmoother(const Operator &oper_,
diag(d),
coeffs(order),
ess_tdof_list(ess_tdofs),
residual(N),
residual(order > 1 ? N : 0),
z(order > 1 ? N : 0),
oper(&oper_) { Setup(); }
#ifdef MFEM_USE_MPI
@@ -375,7 +380,8 @@ OperatorChebyshevSmoother::OperatorChebyshevSmoother(const Operator &oper_,
diag(d),
coeffs(order),
ess_tdof_list(ess_tdofs),
residual(N),
residual(order > 1 ? N : 0),
z(order > 1 ? N : 0),
oper(&oper_)
{
OperatorJacobiSmoother invDiagOperator(diag, ess_tdofs, 1.0);
@@ -422,7 +428,7 @@ void OperatorChebyshevSmoother::Setup()
{
// Invert diagonal
residual.UseDevice(true);
helperVector.UseDevice(true);
z.UseDevice(true);
auto D = diag.Read();
auto X = dinv.Write();
mfem::forall(N, [=] MFEM_HOST_DEVICE (int i) { X[i] = 1.0 / D[i]; });
@@ -432,6 +438,20 @@ void OperatorChebyshevSmoother::Setup()
X[I[i]] = 1.0;
});
const int order_save = order;
order = -1; // avoid early exit in SetOrder() when 'new_order' == 'order'
SetOrder(order_save);
}
void OperatorChebyshevSmoother::SetOrder(int new_order)
{
if (new_order == order) { return; }
order = new_order;
coeffs.SetSize(order);
residual.SetSize(order > 1 ? N : 0);
z.SetSize(order > 1 ? N : 0);
// Set up Chebyshev coefficients
// For reference, see e.g., Parallel multigrid smoothing: polynomial versus
// Gauss-Seidel by Adams et al.
@@ -511,32 +531,55 @@ void OperatorChebyshevSmoother::Mult(const Vector& x, Vector &y) const
MFEM_ABORT("Chebyshev smoother requires operator");
}
residual = x;
helperVector.SetSize(x.Size());
helperVector.UseDevice(true);
y.UseDevice(true);
y = 0.0;
for (int k = 0; k < order; ++k)
// for k = 0, perform:
// r = D^{-1} x
// y = C_0 r
const real_t C_0 = coeffs[0];
auto Dinv = dinv.Read();
auto X = x.Read();
auto Y0 = y.Write();
if (order == 1)
{
// Apply
if (k > 0)
mfem::forall(N, [=] MFEM_HOST_DEVICE (int i)
{
oper->Mult(residual, helperVector);
residual = helperVector;
}
Y0[i] = C_0 * Dinv[i] * X[i];
});
}
else
{
auto R0 = residual.Write();
mfem::forall(N, [=] MFEM_HOST_DEVICE (int i)
{
Y0[i] = C_0 * (R0[i] = Dinv[i] * X[i]);
});
}
// Scale residual by inverse diagonal
const int n = N;
auto Dinv = dinv.Read();
auto R = residual.ReadWrite();
mfem::forall(n, [=] MFEM_HOST_DEVICE (int i) { R[i] *= Dinv[i]; });
for (int k = 1; k < order; ++k)
{
// Apply: z = A r
oper->Mult(residual, z);
// Add weighted contribution to y
// Scale residual by inverse diagonal and add weighted contribution to y:
// r = D^{-1} z
// y += C_k r
const real_t C_k = coeffs[k];
auto Z = z.Read();
auto Y = y.ReadWrite();
auto C = coeffs.Read();
mfem::forall(n, [=] MFEM_HOST_DEVICE (int i) { Y[i] += C[k] * R[i]; });
if (k < order-1)
{
auto R = residual.Write();
mfem::forall(N, [=] MFEM_HOST_DEVICE (int i)
{
Y[i] += C_k * (R[i] = Dinv[i] * Z[i]);
});
}
else
{
mfem::forall(N, [=] MFEM_HOST_DEVICE (int i)
{
Y[i] += C_k * Dinv[i] * Z[i];
});
}
}
}
@@ -3213,7 +3256,7 @@ void ResidualBCMonitor::MonitorResidual(
MPI_Comm comm = iter_solver->GetComm();
if (comm != MPI_COMM_NULL)
{
double glob_bc_norm_squared = 0.0;
real_t glob_bc_norm_squared = 0.0;
MPI_Reduce(&bc_norm_squared, &glob_bc_norm_squared, 1,
MPITypeMap<real_t>::mpi_type,
MPI_SUM, 0, comm);
+9 -8
View File
@@ -380,11 +380,11 @@ public:
void SetPositiveDiagonal(bool pos_diag = true) { use_abs_diag = pos_diag; }
/// Approach the solution of the linear system by applying Jacobi smoothing.
void Mult(const Vector &x, Vector &y) const;
void Mult(const Vector &x, Vector &y) const override;
/** @brief Approach the solution of the transposed linear system by applying
Jacobi smoothing. */
void MultTranspose(const Vector &x, Vector &y) const { Mult(x, y); }
void MultTranspose(const Vector &x, Vector &y) const override { Mult(x, y); }
/** @brief Recompute the diagonal using the method AssembleDiagonal of the
given new Operator, @a op. */
@@ -397,7 +397,7 @@ public:
When the new Operator, @a op, is not a (Par)BilinearForm, any previously
set array of essential true-dofs will be thrown away because in this case
any essential b.c. will be handled by the AssembleDiagonal method. */
void SetOperator(const Operator &op);
void SetOperator(const Operator &op) override;
private:
Vector dinv;
@@ -481,21 +481,22 @@ public:
/** @brief Approach the solution of the linear system by applying Chebyshev
smoothing. */
void Mult(const Vector &x, Vector &y) const;
void Mult(const Vector &x, Vector &y) const override;
/** @brief Approach the solution of the transposed linear system by applying
Chebyshev smoothing. */
void MultTranspose(const Vector &x, Vector &y) const { Mult(x, y); }
void MultTranspose(const Vector &x, Vector &y) const override { Mult(x, y); }
void SetOperator(const Operator &op_)
void SetOperator(const Operator &op_) override
{
oper = &op_;
}
void Setup();
void SetOrder(int new_order);
private:
const int order;
int order;
real_t max_eig_estimate;
const int N;
Vector dinv;
@@ -503,7 +504,7 @@ private:
Array<real_t> coeffs;
const Array<int>& ess_tdof_list;
mutable Vector residual;
mutable Vector helperVector;
mutable Vector z;
const Operator* oper;
};
+1 -42
View File
@@ -14,7 +14,6 @@
#include "../general/forall.hpp"
#include "../general/reducers.hpp"
#include "../general/hash.hpp"
#include "../general/scan.hpp"
#include "vector.hpp"
#ifdef MFEM_USE_OPENMP
@@ -118,7 +117,7 @@ Vector::Vector(const Vector &v)
UseDevice(v.UseDevice());
}
Vector::Vector(Vector &&v) : Vector()
Vector::Vector(Vector &&v)
{
*this = std::move(v);
}
@@ -1253,44 +1252,4 @@ real_t Vector::Sum() const
return res;
}
void Vector::DeleteAt(const Array<int> &indices)
{
const bool use_dev = UseDevice();
Array<int> flag(size);
const auto d_flag = flag.Write(use_dev);
mfem::forall_switch(use_dev, size, [=] MFEM_HOST_DEVICE (int i)
{
d_flag[i] = true;
});
const auto d_indices = indices.Read(use_dev);
mfem::forall_switch(use_dev, indices.Size(), [=] MFEM_HOST_DEVICE (int i)
{
d_flag[d_indices[i]] = false;
});
Array<int> out_idx(size);
auto d_out_idx = out_idx.Write(use_dev);
Array<char> workspace;
// Perform inclusive scan so that the last entry is the new size.
InclusiveScan(use_dev, d_flag, d_out_idx, size, workspace);
Vector copy(*this);
auto d_in = copy.Read(use_dev);
auto d_out = Write(use_dev);
mfem::forall_switch(use_dev, size, [=] MFEM_HOST_DEVICE (int i)
{
if (d_flag[i])
{
// Transform inclusive scan to exclusive by shifting.
const int j = (i > 0) ? d_out_idx[i - 1] : 0;
d_out[j] = d_in[i];
}
});
// Get the new size of the vector. Copy only the last entry.
Memory<int> submem(out_idx.GetMemory(), out_idx.Size() - 1, 1);
size = submem.Read(MemoryClass::HOST, 1)[0];
}
} // namespace mfem
-18
View File
@@ -171,12 +171,6 @@ public:
/// Resize the vector to size @a s using the MemoryType of @a v.
void SetSize(int s, const Vector &v) { SetSize(s, v.GetMemory().GetMemoryType()); }
/// Update \ref Capacity() to @a res (if less than current), keeping existing entries.
void Reserve(int res);
/// Delete entries at @a indices and resize vector accordingly.
void DeleteAt(const Array<int> &indices);
/// Set the Vector data.
/// @warning This method should be called only when OwnsData() is false.
void SetData(real_t *d) { data.Wrap(d, data.Capacity(), false); }
@@ -627,18 +621,6 @@ inline void Vector::SetSize(int s, MemoryType mt)
data.UseDevice(use_dev);
}
inline void Vector::Reserve(int res)
{
if (res > Capacity())
{
Memory<real_t> p(res, data.GetMemoryType());
p.CopyFrom(data, size);
p.UseDevice(data.UseDevice());
data.Delete();
data = p;
}
}
inline void Vector::NewMemoryAndSize(const Memory<real_t> &mem, int s,
bool own_mem)
{
+8 -5
View File
@@ -125,7 +125,8 @@ EXAMPLE_TEST_DIRS := examples
MINIAPP_SUBDIRS = common electromagnetics meshing navier performance tools \
toys nurbs gslib adjoint solvers shifted mtop parelag tribol autodiff dfem \
hooke multidomain dpg hdiv-linear-solver spde diag-smoothers
hooke multidomain dpg hdiv-linear-solver spde diag-smoothers \
benchmarks/ceed-solver-bps
MINIAPP_DIRS := $(addprefix miniapps/,$(MINIAPP_SUBDIRS))
MINIAPP_TEST_DIRS := $(filter-out %/common,$(MINIAPP_DIRS))
MINIAPP_USE_COMMON := $(addprefix miniapps/,electromagnetics meshing tools \
@@ -145,7 +146,7 @@ MFEM_BUILD_DIR ?= .
BUILD_DIR := $(MFEM_BUILD_DIR)
BUILD_REAL_DIR := $(abspath $(BUILD_DIR))
ifneq ($(BUILD_REAL_DIR),$(MFEM_REAL_DIR))
BUILD_SUBDIRS = $(DIRS) config $(EM_DIRS) doc $(TEST_DIRS)
BUILD_SUBDIRS = $(DIRS) config $(EM_DIRS) doc $(TEST_DIRS) tests/benchmarks
CONFIG_FILE_DEF = -DMFEM_CONFIG_FILE='"$(BUILD_REAL_DIR)/config/_config.hpp"'
BLD := $(if $(BUILD_REAL_DIR:$(CURDIR)=),$(BUILD_DIR)/,)
$(if $(word 2,$(BLD)),$(error Spaces in BLD = "$(BLD)" are not supported))
@@ -582,13 +583,14 @@ test-print:
@for dir in $(ALL_TEST_DIRS); do \
$(MAKE) -j1 -C $(BLD)$${dir} test-print; done
ALL_CLEAN_SUBDIRS = $(addsuffix /clean,config $(EM_DIRS) doc $(TEST_DIRS))
ALL_CLEAN_SUBDIRS = $(addsuffix /clean,config $(EM_DIRS) doc $(TEST_DIRS) \
tests/benchmarks)
.PHONY: $(ALL_CLEAN_SUBDIRS) miniapps/clean
miniapps/clean: $(addsuffix /clean,$(MINIAPP_DIRS))
$(ALL_CLEAN_SUBDIRS):
$(MAKE) -C $(BLD)$(@D) $(@F)
clean: $(addsuffix /clean,$(EM_DIRS) $(TEST_DIRS))
clean: $(addsuffix /clean,$(EM_DIRS) $(TEST_DIRS) tests/benchmarks)
rm -f $(addprefix $(BLD),$(foreach d,$(DIRS),$(d)/*.o))
rm -f $(addprefix $(BLD),$(foreach d,$(DIRS),$(d)/*~))
rm -rf $(addprefix $(BLD),*~ libmfem.* deps.mk)
@@ -691,7 +693,8 @@ local-config:
.PHONY: build-config
build-config:
for d in $(BUILD_SUBDIRS); do mkdir -p $(BLD)$${d}; done
for dir in "" $(addsuffix /,config $(EM_DIRS) doc $(TEST_DIRS)); do \
for dir in "" $(addsuffix /,config $(EM_DIRS) doc $(TEST_DIRS) \
tests/benchmarks); do \
printf "# Auto-generated file.\n%s\n%s\n" \
"MFEM_DIR = $(MFEM_REAL_DIR)" \
"include \$$(MFEM_DIR)/$${dir}makefile" \
+28 -167
View File
@@ -32,7 +32,6 @@
#include <cstring>
#include <ctime>
#include <functional>
#include <set>
#include <numeric>
#include <unordered_map>
#include <unordered_set>
@@ -884,6 +883,8 @@ const GeometricFactors* Mesh::GetGeometricFactors(const IntegrationRule& ir,
const int flags,
MemoryType d_mt)
{
MFEM_PERF_FUNCTION;
for (int i = 0; i < geom_factors.Size(); i++)
{
GeometricFactors *gf = geom_factors[i];
@@ -922,103 +923,6 @@ const FaceGeometricFactors* Mesh::GetFaceGeometricFactors(
return gf;
}
const Array<int>& Mesh::GetBdrFaceAttributes() const
{
if (bdr_face_attrs_cache.Size() == 0)
{
std::unordered_map<int, int> f_to_be;
for (int i = 0; i < GetNBE(); ++i)
{
const int f = GetBdrElementFaceIndex(i);
f_to_be[f] = i;
}
const int nf_bdr = GetNFbyType(FaceType::Boundary);
// MFEM_VERIFY(size_t(nf_bdr) == f_to_be.size(), "Incompatible sizes");
bdr_face_attrs_cache.SetSize(nf_bdr);
int f_ind = 0;
const int nf = GetNumFaces();
for (int f = 0; f < nf; ++f)
{
if (!GetFaceInformation(f).IsOfFaceType(FaceType::Boundary))
{
continue;
}
int attribute = -1; // default value
auto iter = f_to_be.find(f);
if (iter != f_to_be.end())
{
const int be = iter->second;
attribute = GetBdrAttribute(be);
}
else
{
// If a boundary face does not correspond to the a boundary element,
// we assign it the default attribute of -1.
}
bdr_face_attrs_cache[f_ind] = attribute;
++f_ind;
}
}
return bdr_face_attrs_cache;
}
const Array<int>& Mesh::GetElementAttributes() const
{
if (elem_attrs_cache.Size() == 0)
{
// re-compute cache
elem_attrs_cache.SetSize(GetNE());
elem_attrs_cache.HostWrite();
for (int i = 0; i < GetNE(); ++i)
{
elem_attrs_cache[i] = GetAttribute(i);
MFEM_ASSERT(elem_attrs_cache[i] > 0,
"Negative attribute on element " << i);
}
}
return elem_attrs_cache;
}
void Mesh::ComputeFaceInfo(FaceType ftype) const
{
auto &fidcs = face_indices[static_cast<int>(ftype)];
auto &ifidcs = inv_face_indices[static_cast<int>(ftype)];
fidcs.SetSize(GetNFbyType(ftype));
fidcs.HostWrite();
ifidcs.reserve(fidcs.Size());
int f_idx = 0;
for (int i = 0; i < GetNumFacesWithGhost(); ++i)
{
const FaceInformation face = GetFaceInformation(i);
if (face.IsNonconformingCoarse() || !face.IsOfFaceType(ftype))
{
continue;
}
fidcs[f_idx] = i;
ifidcs[i] = f_idx;
++f_idx;
}
}
const Array<int> &Mesh::GetFaceIndices(FaceType ftype) const
{
if (face_indices[static_cast<int>(ftype)].Size() == 0)
{
ComputeFaceInfo(ftype);
}
return face_indices[static_cast<int>(ftype)];
}
const std::unordered_map<int, int> &
Mesh::GetInvFaceIndices(FaceType ftype) const
{
if (inv_face_indices[static_cast<int>(ftype)].empty())
{
ComputeFaceInfo(ftype);
}
return inv_face_indices[static_cast<int>(ftype)];
}
void Mesh::DeleteGeometricFactors()
{
for (int i = 0; i < geom_factors.Size(); i++)
@@ -1902,15 +1806,8 @@ void Mesh::Destroy()
TetMemory.Clear();
#endif
elem_attrs_cache.DeleteAll();
bdr_face_attrs_cache.DeleteAll();
attributes.DeleteAll();
bdr_attributes.DeleteAll();
face_indices[0].DeleteAll();
face_indices[1].DeleteAll();
inv_face_indices[0] = std::unordered_map<int, int>();
inv_face_indices[1] = std::unordered_map<int, int>();
}
void Mesh::ResetLazyData()
@@ -1922,48 +1819,36 @@ void Mesh::ResetLazyData()
DeleteGeometricFactors();
nbInteriorFaces = -1;
nbBoundaryFaces = -1;
// set size to 0 so re-computations can potentially avoid a new allocation
bdr_face_attrs_cache.SetSize(0);
elem_attrs_cache.SetSize(0);
}
void Mesh::SetAttributes(bool elem_attrs_changed, bool bdr_face_attrs_changed)
void Mesh::SetAttributes()
{
if (bdr_face_attrs_changed)
Array<int> attribs;
attribs.SetSize(GetNBE());
for (int i = 0; i < attribs.Size(); i++)
{
bdr_face_attrs_cache.SetSize(0); // Invalidate the cache
// Get sorted list of unique boundary element attributes
std::set<int> attribs;
for (int i = 0; i < GetNBE(); i++)
{
attribs.emplace(GetBdrAttribute(i));
}
bdr_attributes.SetSize(attribs.size());
bdr_attributes.HostWrite();
std::copy(attribs.begin(), attribs.end(), bdr_attributes.begin());
if (bdr_attributes.Size() > 0 && bdr_attributes[0] <= 0)
{
MFEM_WARNING("Non-positive attributes on the boundary!");
}
attribs[i] = GetBdrAttribute(i);
}
attribs.Sort();
attribs.Unique();
attribs.Copy(bdr_attributes);
if (bdr_attributes.Size() > 0 && bdr_attributes[0] <= 0)
{
MFEM_WARNING("Non-positive attributes on the boundary!");
}
if (elem_attrs_changed)
attribs.SetSize(GetNE());
for (int i = 0; i < attribs.Size(); i++)
{
// Re-compute the attributes cache
elem_attrs_cache.SetSize(0);
GetElementAttributes();
// Get sorted list of unique element attributes
std::set<int> attribs(elem_attrs_cache.begin(), elem_attrs_cache.end());
attributes.SetSize(attribs.size());
attributes.HostWrite();
std::copy(attribs.begin(), attribs.end(), attributes.begin());
if (attributes.Size() > 0 && attributes[0] <= 0)
{
MFEM_WARNING("Non-positive attributes in the domain!");
}
attribs[i] = GetAttribute(i);
}
attribs.Sort();
attribs.Unique();
attribs.Copy(attributes);
if (attributes.Size() > 0 && attributes[0] <= 0)
{
MFEM_WARNING("Non-positive attributes in the domain!");
}
}
@@ -4571,10 +4456,6 @@ Mesh::Mesh(const Mesh &mesh, bool copy_nodes)
Nodes = mesh.Nodes;
own_nodes = 0;
}
// copy attribute caches
elem_attrs_cache = mesh.elem_attrs_cache;
bdr_face_attrs_cache = mesh.bdr_face_attrs_cache;
}
Mesh::Mesh(Mesh &&mesh) : Mesh()
@@ -7842,12 +7723,6 @@ void Mesh::GetBdrElementAdjacentElement2(
void Mesh::SetAttribute(int i, int attr)
{
elements[i]->SetAttribute(attr);
if (elem_attrs_cache.Size() == GetNE())
{
// update the existing cache instead of deleting it
elem_attrs_cache.HostReadWrite();
elem_attrs_cache[i] = attr;
}
if (ncmesh) ncmesh->SetAttribute(i, attr);
}
@@ -8182,12 +8057,6 @@ void Mesh::GenerateFaces()
FreeElement(f);
}
// delete caches
face_indices[0].SetSize(0);
face_indices[1].SetSize(0);
inv_face_indices[0].clear();
inv_face_indices[1].clear();
// (re)generate the interior faces and the info for them
faces.SetSize(nfaces);
faces_info.SetSize(nfaces);
@@ -10983,15 +10852,6 @@ void Mesh::Swap(Mesh& other, bool non_geometry)
mfem::Swap(nodes_sequence, other.nodes_sequence);
mfem::Swap(last_operation, other.last_operation);
}
// copy attribute caches
mfem::Swap(elem_attrs_cache, other.elem_attrs_cache);
mfem::Swap(bdr_face_attrs_cache, other.bdr_face_attrs_cache);
mfem::Swap(face_indices[0], other.face_indices[0]);
mfem::Swap(face_indices[1], other.face_indices[1]);
inv_face_indices[0].swap(other.inv_face_indices[0]);
inv_face_indices[1].swap(other.inv_face_indices[1]);
}
void Mesh::GetElementData(const Array<Element*> &elem_array, int geom,
@@ -13399,7 +13259,7 @@ void Mesh::ScaleElements(real_t sf)
delete [] vn;
}
void Mesh::Transform(std::function<void(const Vector &, Vector&)> f)
void Mesh::Transform(void (*f)(const Vector&, Vector&))
{
// TODO: support for different new spaceDim.
if (Nodes == NULL)
@@ -13412,7 +13272,7 @@ void Mesh::Transform(std::function<void(const Vector &, Vector&)> f)
vold(j) = vertices[i](j);
}
vnew.SetData(vertices[i]());
f(vold, vnew);
(*f)(vold, vnew);
}
}
else
@@ -14791,6 +14651,7 @@ GeometricFactors::GeometricFactors(const GridFunction &nodes,
void GeometricFactors::Compute(const GridFunction &nodes,
MemoryType d_mt)
{
MFEM_PERF_FUNCTION;
const FiniteElementSpace *fespace = nodes.FESpace();
const FiniteElement *fe = fespace->GetTypicalFE();
+11 -61
View File
@@ -107,11 +107,6 @@ protected:
Array<Element *> boundary;
Array<Element *> faces;
/// internal cache for element attributes
mutable Array<int> elem_attrs_cache;
/// internal cache for boundary element attributes
mutable Array<int> bdr_face_attrs_cache;
/** @brief This structure stores the low level information necessary to
interpret the configuration of elements on a specific face. This
information can be accessed using methods like GetFaceElements(),
@@ -278,13 +273,6 @@ protected:
// used during NC mesh initialization only
Array<Triple<int, int, int> > tmp_vertex_parents;
/// cache for FaceIndices(ftype)
mutable Array<int> face_indices[2];
/// cache for FaceIndices(ftype)
mutable std::unordered_map<int, int> inv_face_indices[2];
/// compute face_indices[ftype] and inv_face_indices[type]
void ComputeFaceInfo(FaceType ftype) const;
public:
typedef Geometry::Constants<Geometry::SEGMENT> seg_t;
@@ -319,11 +307,6 @@ public:
// (true) is set in mesh_readers.cpp.
static bool remove_unused_vertices;
/// Map from boundary or interior face indices to mesh face indices.
const Array<int>& GetFaceIndices(FaceType ftype) const;
/// Inverse of the map FaceIndices(ftype)
const std::unordered_map<int, int>& GetInvFaceIndices(FaceType ftype) const;
protected:
Operation last_operation;
@@ -605,10 +588,9 @@ protected:
void Loader(std::istream &input, int generate_edges = 0,
std::string parse_tag = "");
/** @brief If NURBS mesh, write NURBS format. If NCMesh, write mfem v1.1
format. If section_delimiter is empty, write mfem v1.0 format. Otherwise,
write mfem v1.2 format with the given section_delimiter at the end.
/** If NURBS mesh, write NURBS format. If NCMesh, write mfem v1.1 format.
If section_delimiter is empty, write mfem v1.0 format. Otherwise, write
mfem v1.2 format with the given section_delimiter at the end.
If @a comments is non-empty, it will be printed after the first line of
the file, and each line should begin with '#'. */
void Printer(std::ostream &os = mfem::out,
@@ -1139,14 +1121,13 @@ public:
Mesh vertices or nodes are set. */
virtual void Finalize(bool refine = false, bool fix_orientation = false);
/// @brief Determine the sets of unique attribute values in domain if @a
/// elem_attrs_changed and boundary elements if @a bdr_face_attrs_changed.
/// @brief Determine the sets of unique attribute values in domain and
/// boundary elements.
///
/// Separately scan the domain and boundary elements to generate unique,
/// sorted sets of the element attribute values present in the mesh and
/// store these in the Mesh::attributes and Mesh::bdr_attributes arrays.
virtual void SetAttributes(bool elem_attrs_changed = true,
bool bdr_face_attrs_changed = true);
virtual void SetAttributes();
/// Check (and optionally attempt to fix) the orientation of the elements
/** @param[in] fix_it If `true`, attempt to fix the orientations of some
@@ -2272,7 +2253,7 @@ public:
void ScaleSubdomains (real_t sf);
void ScaleElements (real_t sf);
void Transform(std::function<void(const Vector &, Vector&)> f);
void Transform(void (*f)(const Vector&, Vector&));
void Transform(VectorCoefficient &deformation);
/** @brief This function should be called after the mesh node coordinates
@@ -2285,35 +2266,6 @@ public:
method does not modify the nodes. */
void NodesUpdated() { DeleteGeometricFactors(); }
/// @brief Returns the attributes for all elements in this mesh. The i'th
/// entry of the array is the attribute of the i'th element of the mesh.
///
/// The returned array points to an internal object that may be invalidated
/// by mesh operations such as refinement or any element attributes are
/// modified. Since not all such modifications can be tracked by the Mesh
/// class (e.g. if a user calls GetElement() then changes the element
/// attribute directly), one needs to account for such changes by calling the
/// method SetAttributes().
const Array<int>& GetElementAttributes() const;
/// @brief Returns the attributes for all boundary elements in this mesh.
///
/// The face restriction will give "face E-vectors" on the boundary that
/// are numbered in the order of the faces of mesh. This numbering will be
/// different than the numbering of the boundary elements. We compute
/// mappings so that the array `bdr_attributes[i]` gives the boundary
/// attribute of the `i`th boundary face in the mesh face order.
/// Attributes <= 0 indicate there is no boundary element and should be
/// skipped.
///
/// The returned array points to an internal object that may be invalidated
/// by mesh operations such as refinement or any element attributes are
/// modified. Since not all such modifications can be tracked by the Mesh
/// class (e.g. if a user calls GetElement() then changes the element
/// attribute directly), one needs to account for such changes by calling the
/// method SetAttributes().
const Array<int>& GetBdrFaceAttributes() const;
/// @}
/// @anchor mfem_Mesh_gf_nodes
@@ -2530,12 +2482,10 @@ public:
/// Print the mesh to the given stream using Netgen/Truegrid format.
virtual void PrintXG(std::ostream &os = mfem::out) const;
/** @brief Print the mesh to the given stream using the default MFEM mesh
format.
\see mfem::ofgzstream() for on-the-fly compression of ascii outputs. If
@a comments is non-empty, it will be printed after the first line of the
file, and each line should begin with '#'. */
/// Print the mesh to the given stream using the default MFEM mesh format.
/// \see mfem::ofgzstream() for on-the-fly compression of ascii outputs. If
/// @a comments is non-empty, it will be printed after the first line of the
/// file, and each line should begin with '#'.
virtual void Print(std::ostream &os = mfem::out,
const std::string &comments = "") const
{ Printer(os, "", comments); }
+21 -85
View File
@@ -53,7 +53,7 @@ KnotVector::KnotVector(int order, int NCP)
}
KnotVector::KnotVector(int order, const Vector& intervals,
const Array<int>& continuity)
const Array<int>& continuity )
{
// NOTE: This may need to be generalized to support periodicity
// in the future.
@@ -151,7 +151,7 @@ void KnotVector::UniformRefinement(Vector &newknots, int rf) const
{
for (int m = 1; m < rf; ++m)
{
newknots(j) = ((1.0 - (m * h)) * knot(i)) + (m * h * knot(i+1));
newknots(j) = m * h * (knot(i) + knot(i+1));
j++;
}
}
@@ -340,7 +340,7 @@ void KnotVector::PrintFunctions(std::ostream &os, int samples) const
}
}
// Routine from "The NURBS Book" - 2nd ed - Piegl and Tiller
// Routine from "The NURBS book" - 2nd ed - Piegl and Tiller
// Algorithm A2.2 p. 70
void KnotVector::CalcShape(Vector &shape, int i, real_t xi) const
{
@@ -367,7 +367,7 @@ void KnotVector::CalcShape(Vector &shape, int i, real_t xi) const
}
}
// Routine from "The NURBS Book" - 2nd ed - Piegl and Tiller
// Routine from "The NURBS book" - 2nd ed - Piegl and Tiller
// Algorithm A2.3 p. 72
void KnotVector::CalcDShape(Vector &grad, int i, real_t xi) const
{
@@ -425,7 +425,7 @@ void KnotVector::CalcDShape(Vector &grad, int i, real_t xi) const
}
}
// Routine from "The NURBS Book" - 2nd ed - Piegl and Tiller
// Routine from "The NURBS book" - 2nd ed - Piegl and Tiller
// Algorithm A2.3 p. 72
void KnotVector::CalcDnShape(Vector &gradn, int n, int i, real_t xi) const
{
@@ -545,11 +545,11 @@ void KnotVector::FindMaxima(Array<int> &ks, Vector &xi, Vector &u) const
int i = j - d;
if (isElement(i))
{
arg1 = std::numeric_limits<real_t>::epsilon() / 2_r;
arg1 = 1e-16;
CalcShape(shape, i, arg1);
max1 = shape[d];
arg2 = 1_r - arg1;
arg2 = 1-(1e-16);
CalcShape(shape, i, arg2);
max2 = shape[d];
@@ -587,9 +587,9 @@ void KnotVector::FindMaxima(Array<int> &ks, Vector &xi, Vector &u) const
}
}
// Routine from "The NURBS Book" - 2nd ed - Piegl and Tiller
// Routine from "The NURBS book" - 2nd ed - Piegl and Tiller
// Algorithm A9.1 p. 369
void KnotVector::FindInterpolant(Array<Vector*> &x, bool reuse_inverse)
void KnotVector::FindInterpolant(Array<Vector*> &x)
{
int order = GetOrder();
int ncp = GetNCP();
@@ -597,93 +597,29 @@ void KnotVector::FindInterpolant(Array<Vector*> &x, bool reuse_inverse)
// Find interpolation points
Vector xi_args, u_args;
Array<int> i_args;
FindMaxima(i_args, xi_args, u_args);
FindMaxima(i_args,xi_args, u_args);
// Assemble collocation matrix
#ifdef MFEM_USE_LAPACK
// If using LAPACK, we use banded matrix storage (order + 1 nonzeros per row).
// Find banded structure of matrix.
int KL = 0; // Number of subdiagonals
int KU = 0; // Number of superdiagonals
Vector shape(order+1);
DenseMatrix A(ncp,ncp);
A = 0.0;
for (int i = 0; i < ncp; i++)
{
CalcShape(shape, i_args[i], xi_args[i]);
for (int p = 0; p < order+1; p++)
{
const int col = i_args[i] + p;
if (col < i)
{
KL = std::max(KL, i - col);
}
else if (i < col)
{
KU = std::max(KU, col - i);
}
A(i,i_args[i] + p) = shape[p];
}
}
const int LDAB = (2*KL) + KU + 1;
const int N = ncp;
fact_AB.SetSize(LDAB, N);
#else
// Without LAPACK, we store and invert a DenseMatrix (inefficient).
if (!reuse_inverse)
{
A_coll_inv.SetSize(ncp, ncp);
A_coll_inv = 0.0;
}
#endif
Vector shape(order+1);
if (!reuse_inverse) // Set collocation matrix entries
{
for (int i = 0; i < ncp; i++)
{
CalcShape(shape, i_args[i], xi_args[i]);
for (int p = 0; p < order+1; p++)
{
const int j = i_args[i] + p;
#ifdef MFEM_USE_LAPACK
fact_AB(KL+KU+i-j,j) = shape[p];
#else
A_coll_inv(i,j) = shape[p];
#endif
}
}
}
// Solve the system
#ifdef MFEM_USE_LAPACK
const int NRHS = x.Size();
DenseMatrix B(N, NRHS);
for (int j=0; j<NRHS; ++j)
{
for (int i=0; i<N; ++i) { B(i, j) = (*x[j])[i]; }
}
if (reuse_inverse)
{
BandedFactorizedSolve(KL, KU, fact_AB, B, false, fact_ipiv);
}
else
{
BandedSolve(KL, KU, fact_AB, B, fact_ipiv);
}
for (int j=0; j<NRHS; ++j)
{
for (int i=0; i<N; ++i) { (*x[j])[i] = B(i, j); }
}
#else
if (!reuse_inverse) { A_coll_inv.Invert(); }
// Solve problems
A.Invert();
Vector tmp;
for (int i = 0; i < x.Size(); i++)
for (int i= 0; i < x.Size(); i++)
{
tmp = *x[i];
A_coll_inv.Mult(tmp, *x[i]);
A.Mult(tmp,*x[i]);
}
#endif
}
int KnotVector::findKnotSpan(real_t u) const
@@ -1485,7 +1421,7 @@ void NURBSPatch::DegreeElevate(int t)
}
}
// Routine from "The NURBS Book" - 2nd ed - Piegl and Tiller
// Routine from "The NURBS book" - 2nd ed - Piegl and Tiller
void NURBSPatch::DegreeElevate(int dir, int t)
{
if (dir >= kv.Size() || dir < 0)
@@ -2449,7 +2385,7 @@ NURBSExtension::NURBSExtension(Mesh *mesh_array[], int num_pieces)
}
NURBSExtension::NURBSExtension(const Mesh *patch_topology,
const Array<const NURBSPatch*> &patches_)
const Array<const NURBSPatch*> patches_)
{
// Basic topology checks
MFEM_VERIFY(patches_.Size() > 0, "Must have at least one patch");
+4 -18
View File
@@ -143,13 +143,8 @@ public:
/** @brief Global curve interpolation through the points @a x (overwritten).
@a x is an array with the length of the spatial dimension containing
vectors with spatial coordinates. The control points of the interpolated
curve are returned in @a x in the same form.
The inverse of the collocation matrix, used in the interpolation, is
stored for repeated calls and used if @a reuse_inverse is true. Reuse is
valid only if this KnotVector has not changed since the initial call with
@a reuse_inverse false. */
void FindInterpolant(Array<Vector*> &x, bool reuse_inverse = false);
curve are returned in @a x in the same form. */
void FindInterpolant(Array<Vector*> &x);
/** Set @a diff, comprised of knots in @a kv not contained in this KnotVector.
@a kv must be of the same order as this KnotVector. The current
@@ -207,14 +202,6 @@ public:
/** Flag to indicate whether the KnotVector has been coarsened, which means
it is ready for non-nested refinement. */
bool coarse;
#ifdef MFEM_USE_LAPACK
// Data for reusing banded matrix factorization in FindInterpolant().
DenseMatrix fact_AB; /// Banded matrix factorization
Array<int> fact_ipiv; /// Row pivot indices
#else
DenseMatrix A_coll_inv; /// Collocation matrix inverse
#endif
};
@@ -298,7 +285,7 @@ public:
includes the weight. The array of control point coordinates stores each
point's coordinates contiguously, and points are ordered in a standard
ijk grid ordering. */
NURBSPatch(Array<const KnotVector *> &kv_, int dim_,
NURBSPatch(Array<const KnotVector *> &kv_, int dim_,
const real_t* control_points);
/// Constructor for a patch of dimension equal to the size of @a kv.
@@ -713,8 +700,7 @@ public:
NURBSExtension(Mesh *mesh_array[], int num_pieces);
NURBSExtension(const Mesh *patch_topology,
const Array<const NURBSPatch*> &patches_);
NURBSExtension(const Mesh *patch_topology, const Array<const NURBSPatch*> p);
/// Copy assignment not supported.
NURBSExtension& operator=(const NURBSExtension&) = delete;
+8 -14
View File
@@ -1590,27 +1590,21 @@ void ParMesh::DistributeAttributes(Array<int> &attr)
delete [] glb_attr_marker;
}
void ParMesh::SetAttributes(bool elem_attrs_changed, bool bdr_attrs_changed)
void ParMesh::SetAttributes()
{
// Determine the attributes occurring in local interior and boundary elements
Mesh::SetAttributes(elem_attrs_changed, bdr_attrs_changed);
Mesh::SetAttributes();
if (bdr_attrs_changed)
DistributeAttributes(bdr_attributes);
if (bdr_attributes.Size() > 0 && bdr_attributes[0] <= 0)
{
DistributeAttributes(bdr_attributes);
if (bdr_attributes.Size() > 0 && bdr_attributes[0] <= 0)
{
MFEM_WARNING("Non-positive boundary element attributes found!");
}
MFEM_WARNING("Non-positive boundary element attributes found!");
}
if (elem_attrs_changed)
DistributeAttributes(attributes);
if (attributes.Size() > 0 && attributes[0] <= 0)
{
DistributeAttributes(attributes);
if (attributes.Size() > 0 && attributes[0] <= 0)
{
MFEM_WARNING("Non-positive element attributes found!");
}
MFEM_WARNING("Non-positive element attributes found!");
}
}
+1 -2
View File
@@ -394,8 +394,7 @@ public:
void Finalize(bool refine = false, bool fix_orientation = false) override;
void SetAttributes(bool elem_attrs_changed = true,
bool bdr_attrs_changed = true) override;
void SetAttributes() override;
/// Checks if any rank in the mesh has boundary elements
bool HasBoundaryElements() const override;
+9 -1
View File
@@ -257,7 +257,15 @@ template <typename SubMeshT>
void AddBoundaryElements(SubMeshT &mesh,
const std::unordered_map<int,int> &lface_to_boundary_attribute)
{
const int num_codim_1 = mesh.GetNumFaces();
mesh.Dimension();
const int num_codim_1 = [&mesh]()
{
auto Dim = mesh.Dimension();
if (Dim == 1) { return mesh.GetNV(); }
else if (Dim == 2) { return mesh.GetNEdges(); }
else if (Dim == 3) { return mesh.GetNFaces(); }
else { MFEM_ABORT("Invalid dimension."); return -1; }
}();
if (mesh.Dimension() == 3)
{
@@ -0,0 +1,156 @@
// Copyright (c) 2010-2025, Lawrence Livermore National Security, LLC. Produced
// at the Lawrence Livermore National Laboratory. All Rights reserved. See files
// LICENSE and NOTICE for details. LLNL-CODE-806117.
//
// This file is part of the MFEM library. For more information and source code
// availability visit https://mfem.org.
//
// MFEM is free software; you can redistribute it and/or modify it under the
// terms of the BSD-3 license. We welcome feedback and contributions, see file
// CONTRIBUTING.md for details.
#ifndef __KERSHAW_HPP__
#define __KERSHAW_HPP__
#include "mfem.hpp"
namespace mfem
{
// 1D transformation at the right boundary.
real_t right(const real_t eps, const real_t x)
{
return (x <= 0.5) ? (2-eps) * x : 1 + eps*(x-1);
}
// 1D transformation at the left boundary
real_t left(const real_t eps, const real_t x)
{
return 1-right(eps,1-x);
}
// Transition from a value of "a" for x=0, to a value of "b" for x=1. Smoothness
// is controlled by the parameter "s", taking values 0, 1, or 2.
real_t step(const real_t a, const real_t b, real_t x, int s)
{
if (x <= 0) { return a; }
if (x >= 1) { return b; }
switch (s)
{
case 0:
default:
return a + (b-a) * (x);
case 1: return a + (b-a) * (x*x*(3-2*x));
case 2: return a + (b-a) * (x*x*x*(x*(6*x-15)+10));
}
}
// 3D version of a generalized Kershaw mesh transformation, see D. Kershaw,
// "Differencing of the diffusion equation in Lagrangian hydrodynamic codes",
// JCP, 39:375395, 1981.
//
// The input mesh should be Cartesian nx x ny x nz with nx divisible by 6 and
// ny, nz divisible by 2.
//
// The eps parameters are in (0, 1]. Uniform mesh is recovered for epsy=epsz=1.
void kershaw(const real_t epsy, const real_t epsz, const int smoothness,
const real_t x, const real_t y, const real_t z,
real_t &X, real_t &Y, real_t &Z)
{
X = x;
int layer = x*6.0;
real_t lambda = (x-layer/6.0)*6;
// The x-range is split in 6 layers going from left-to-left, left-to-right,
// right-to-left (2 layers), left-to-right and right-to-right yz-faces.
switch (layer)
{
case 0:
Y = left(epsy, y);
Z = left(epsz, z);
break;
case 1:
case 4:
Y = step(left(epsy, y), right(epsy, y), lambda, smoothness);
Z = step(left(epsz, z), right(epsz, z), lambda, smoothness);
break;
case 2:
Y = step(right(epsy, y), left(epsy, y), lambda/2, smoothness);
Z = step(right(epsz, z), left(epsz, z), lambda/2, smoothness);
break;
case 3:
Y = step(right(epsy, y), left(epsy, y), (1+lambda)/2, smoothness);
Z = step(right(epsz, z), left(epsz, z), (1+lambda)/2, smoothness);
break;
default:
Y = right(epsy, y);
Z = right(epsz, z);
break;
}
}
struct KershawTransformation : VectorCoefficient
{
real_t epsy, epsz;
int dim, s;
KershawTransformation(int dim_, real_t epsy_, real_t epsz_, int s_=0)
: VectorCoefficient(dim_), epsy(epsy_), epsz(epsz_), dim(dim_), s(s_) { }
using VectorCoefficient::Eval;
void Eval(Vector &V, ElementTransformation &T,
const IntegrationPoint &ip) override
{
real_t xyz[3];
Vector transip(xyz, 3);
T.Transform(ip, transip);
if (dim == 1)
{
V[0] = xyz[0]; // no transformation in 1D
}
else if (dim == 2)
{
real_t z=0, zt;
kershaw(epsy, epsz, s, xyz[0], xyz[1], z, V[0], V[1], zt);
}
else // dim == 3
{
kershaw(epsy, epsz, s, xyz[0], xyz[1], xyz[2], V[0], V[1], V[2]);
}
}
};
ParMesh CreateKershawMesh(int nx, int ny, int nz, real_t epsy, real_t epsz)
{
const bool sfc_order = true;
Mesh serial_mesh;
if (nx > 0 && ny == 0 && nz == 0)
{
serial_mesh = Mesh::MakeCartesian1D(nx, 1.0);
}
else if (nx > 0 && ny > 0 && nz == 0)
{
serial_mesh = Mesh::MakeCartesian2D(nx, ny, Element::QUADRILATERAL,
false, 1, 1, sfc_order);
}
else if (nx > 0 && ny > 0 && nz > 0)
{
serial_mesh = Mesh::MakeCartesian3D(nx, ny, nz, Element::HEXAHEDRON,
1, 1, 1, sfc_order);
}
else
{
MFEM_ABORT("Bad grid size");
}
KershawTransformation kt(serial_mesh.Dimension(), epsy, epsz);
serial_mesh.Transform(kt);
return ParMesh(MPI_COMM_WORLD, serial_mesh);
}
ParMesh CreateKershawMesh(int n, real_t eps)
{
return CreateKershawMesh(n, n, n, eps, eps);
}
}
#endif
@@ -0,0 +1,77 @@
# Copyright (c) 2010-2025, Lawrence Livermore National Security, LLC. Produced
# at the Lawrence Livermore National Laboratory. All Rights reserved. See files
# LICENSE and NOTICE for details. LLNL-CODE-806117.
#
# This file is part of the MFEM library. For more information and source code
# availability visit https://mfem.org.
#
# MFEM is free software; you can redistribute it and/or modify it under the
# terms of the BSD-3 license. We welcome feedback and contributions, see file
# CONTRIBUTING.md for details.
# Use the MFEM build directory
MFEM_DIR ?= ../../..
MFEM_BUILD_DIR ?= ../../..
MFEM_INSTALL_DIR ?= ../../../mfem
SRC = $(if $(MFEM_DIR:../../..=),$(MFEM_DIR)/miniapps/benchmarks/ceed-solver-bps/,)
CONFIG_MK = $(or $(wildcard $(MFEM_BUILD_DIR)/config/config.mk),\
$(wildcard $(MFEM_INSTALL_DIR)/share/mfem/config.mk))
MFEM_LIB_FILE = mfem_is_not_built
-include $(CONFIG_MK)
SEQ_MINIAPPS =
PAR_MINIAPPS = solver-bp
ifeq ($(MFEM_USE_MPI),NO)
MINIAPPS = $(SEQ_MINIAPPS)
else
MINIAPPS = $(PAR_MINIAPPS) $(SEQ_MINIAPPS)
endif
EXTRA_SOURCES = preconditioners.cpp
EXTRA_HEADERS = kershaw.hpp rhs.hpp preconditioners.hpp
EXTRA_OBJECTS = $(EXTRA_SOURCES:.cpp=.o)
.SUFFIXES:
.SUFFIXES: .o .cpp .mk
.PHONY: all clean clean-build clean-exec
.PRECIOUS: %.o
# Remove built-in rules
%: %.cpp
%.o: %.cpp
all: $(MINIAPPS)
# Rule for building solver-bp
solver-bp: solver-bp.o $(addprefix $(SRC),$(EXTRA_HEADERS)) \
$(EXTRA_OBJECTS) $(MFEM_LIB_FILE) $(CONFIG_MK)
$(MFEM_CXX) $(MFEM_LINK_FLAGS) $< -o $@ $(EXTRA_OBJECTS) $(MFEM_LIBS)
# Rules for compiling *.o files
# -I$(MFEM_DIR) is needed for "general/forall.hpp" for out-of-source builds
%.o: $(SRC)%.cpp $(wildcard $(SRC)%.hpp) $(CONFIG_MK)
$(MFEM_CXX) $(MFEM_FLAGS) -I$(MFEM_DIR) -c $(<) -o $(@)
MFEM_TESTS = MINIAPPS
include $(MFEM_TEST_MK)
# Testing: Specific execution options
RUN_MPI = $(MFEM_MPIEXEC) $(MFEM_MPIEXEC_NP) $(MFEM_MPI_NP)
solver-bp-test-par: solver-bp
@$(call mfem-test,$<, $(RUN_MPI), CEED Solver BP,,SKIP-NO-VIS)
# Testing: "test" target and mfem-test* variables are defined in config/test.mk
# Generate an error message if the MFEM library is not built and exit
$(MFEM_LIB_FILE):
$(error The MFEM library is not built)
clean: clean-build clean-exec
clean-build:
rm -f *.o *~ $(SEQ_MINIAPPS) $(PAR_MINIAPPS) $(EXTRA_OBJECTS)
rm -rf *.dSYM *.TVD.*breakpoints
clean-exec:
@true
@@ -0,0 +1,129 @@
import csv
from pylab import *
fields=[
['code ID', 'str'],
['preconditioner ID', 'str'],
['machine ID', 'str'],
['number of nodes', 'int'],
['number of MPI ranks', 'int'],
['n_x', 'int'], ['n_y', 'int'], ['n_z', 'int'],
['solution polynomial degree', 'int'],
['number of 1D quadrature points', 'float'],
['eps_y', 'float'], ['eps_z', 'float'],
['ndofs (including Dirichlet boundary)', 'int'],
['niter', 'int'],
['initial residual', 'float'], ['final residual', 'float'],
['error', 'float'],
['t_setup (preconditioner setup)', 'float'],
['t_solve (total iter time)', 'float']]
fields_dict=dict(fields)
def convert(obj, type_str):
ctor=getattr(__builtins__, type_str)
return ctor(obj)
input_csv='run-001.csv'
print('reading %s ...' % input_csv)
runs = []
with open(input_csv) as csvfile:
csvreader = csv.DictReader(csvfile, fieldnames=[f[0] for f in fields],
restkey='additional notes')
for row in csvreader:
for i in fields_dict:
row[i]=convert(row[i], fields_dict[i])
runs.append(row)
orders=[r['solution polynomial degree'] for r in runs]
orders=unique(orders) # numpy function
# orders=[1]
nps=[r['number of MPI ranks'] for r in runs]
nps=unique(nps)
if len(nps) > 1:
print('multiple num-ranks present: %s' % nps)
quit()
np=nps[0]
# plot fx (or fx/fn) vs fy, (or fx/fn/fy, etc) for all orders
fn='number of MPI ranks'
fx='ndofs (including Dirichlet boundary)'
fy='t_solve (total iter time)'
# fy='niter'
# fy='error'
fz='niter'
figure()
for p in orders:
rr=[r for r in runs if (r['solution polynomial degree']==p and
r['niter']>0)]
if len(rr)==0:
continue
# pl_data=asarray([[r[fx],r[fx]/r[fy]] for r in rr])
# pl_data=asarray([[r[fx],r[fy]] for r in rr])
# pl_data=asarray([[r[fx],r[fx]/(r[fy]/r[fz])] for r in rr])
pl_data=asarray([[r[fx]/r[fn],r[fx]/r[fn]/r[fy]] for r in rr])
# pl_data=asarray([[r[fx]/r[fn],r[fy]] for r in rr])
plot(pl_data[:,0],pl_data[:,1], 'o-', label='p=%i'%p)
rnx=asarray([r['n_x'] for r in rr])
rerr=asarray([r['error'] for r in rr])
rate=arange(1.0,len(rnx))
for l in range(1,len(rnx)):
rate[l-1]=log(rerr[l-1]/rerr[l])/log(rnx[l]/rnx[l-1])
set_printoptions(formatter={'float':"{:6.2f}".format},linewidth=120)
print(f"p={p} rate:{rate}")
# xscale('log', basex=10) # older matplotlib
xscale('log', base=10)
# xlim(4e4,3.1e7)
xlim(4e4,5e6)
# yscale('log', basey=10) # older matplotlib
# yscale('log', base=10)
# ylim(1e5,2e7)
# ylim(0,2.55e7)
# ylim(0,3.25e7)
# ylim(0,5e6)
ymin,ymax=ylim()
ylim(0,ymax)
# ylim(1e-2,2e1)
# ylim(3e-3,6e-2)
# xlabel(fx)
# xlabel('# DOFs')
xlabel('# DOFs / # Ranks')
# ylabel(fx + ' / ' + fy)
# ylabel(fy)
# ylabel('# DOFs / t_solve')
ylabel('# DOFs / # Ranks / t_solve')
# ylabel('t_solve')
# ylabel('# DOFs / (t_solve / # Iter)')
# ylabel('# Iter')
# ylabel('L2 error')
# ylabel('Grad L2 error')
grid('on', color='gray', ls='dotted')
grid('on', axis='both', which='minor', color='gray', ls='dotted')
legend(ncol=2, loc='best')
ranks='1 MPI rank'
if np > 1:
ranks='%s MPI ranks' % (np,np)
hypre='hypre CPU'
# hypre='hypre HIP'
# prec=hypre+', p-MG(1,1)'
prec=hypre+', LOR'
# prec='Jacobi'
# eps='1'
eps='0.3'
mfem='MFEM CPU'
# mfem='MFEM HIP'
title(mfem + ', ' + prec + ', $\\varepsilon = ' + eps + '$, ' + ranks)
if 1: # write .pdf file?
pdf_file='plot.pdf'
print('saving figure --> %s'%pdf_file)
savefig(pdf_file, format='pdf', bbox_inches='tight')
if 0: # show the figures?
print('\nshowing figures ...')
show()

Some files were not shown because too many files have changed in this diff Show More