Compare commits

..
Author SHA1 Message Date
camierjs f2a424720b Merge branch 'main' into hpcftools/flowSolver-gpu 2025-06-30 11:17:56 -07:00
camierjs e5fad8067c Merge branch 'master' into hpcftools/flowSolver-gpu 2025-06-30 11:17:45 -07:00
camierjs 0b5bb52996 miniapps/navier/incompressible_navier_dfem 2025-06-30 11:17:11 -07:00
camierjs 563bfe9514 Merge branch 'master' 2025-06-30 08:43:14 -07:00
camierjs 9afef578d5 Merge branch 'master' into hpcftools/flowSolver-gpu 2025-05-02 09:59:44 -07:00
camierjs 9ecc414c62 dot reduced tests 2024-12-04 18:25:05 -08:00
camierjs 263b9d32c1 Merge branch 'hpcftools/flowSolver' 2024-12-04 11:59:05 -08:00
Mathias Rainer Schmidt 74e4ad3e2c - updated flow solver
- added comments to Blf and Lf contributions
- split setup and step into vel, auxiliary and pressure part
2024-11-25 11:04:30 -08:00
camierjs 79039e0f6f Switched to PA 2024-11-21 10:52:39 -08:00
Mathias Rainer Schmidt cf9fcd8dde Merge remote-tracking branch 'origin/master' into hpcftools/flowSolver 2024-11-14 13:52:27 -08:00
camierjs c857fde13b Merge branch 'hpcftools/flowSolver' 2024-11-01 15:59:50 -07:00
Mathias Rainer Schmidt fa8617ada3 - added partial assembly option 2024-11-01 13:20:29 -07:00
camierjs c06cbb69d5 Setup and cleanup 2024-10-30 11:37:00 -07:00
Mathias Rainer Schmidt f7e5db2cea - added ortho solver to phi field 2024-10-24 16:06:37 -07:00
Mathias Rainer Schmidt 60c11776b6 - added executable 2024-10-21 15:57:23 -07:00
Mathias Rainer Schmidt 6238f8ca76 - update solution 2024-10-16 16:00:08 -07:00
Mathias Rainer Schmidt 6a256db9aa - added linear solvers to step 2024-10-16 15:56:46 -07:00
Mathias Rainer Schmidt e0fb9658ca - added linear form integrators 2024-10-15 17:30:09 -07:00
Mathias Rainer Schmidt a1089efac3 - added BilinearForms 2024-10-15 13:46:36 -07:00
Mathias Rainer Schmidt 83f7f769dc - inital flow solver commit 2024-10-15 12:44:25 -07:00
193 changed files with 4049 additions and 8112 deletions
-154
View File
@@ -1,154 +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.
---
name: Sanitizer Config
description: Sets up environment variables for MFEM sanitizer workflow
inputs:
DEBUG:
description: If true, use intermediate caches to speed up the workflow
by reusing previous builds.
default: false
REPOSITORY:
description: Repository to checkout
default: mfem/mfem
BRANCH:
description: Branch to checkout
default: ubsan
CLANG_VER:
description: CLANG version to use
default: 18
# https://github.com/llvm/llvm-project/releases
LLVM_VER:
description: LLVM version to use
default: 19.1.7
# https://github.com/hypre-space/hypre/releases
HYPRE_VER:
description: HYPRE version to use
default: 2.19.0
METIS_VER:
description: METIS version to use
default: 4.0.3
CTEST:
description: CTest command to use
default: ctest -j --test-load $(nproc)
--schedule-random
--stop-on-failure --output-on-failure
--test-dir
# https://clang.llvm.org/docs/AddressSanitizer.html
ASAN_OPTIONS:
default: detect_leaks=1,
strict_init_order=1,
strict_string_checks=1,
check_initialization_order=1,
detect_stack_use_after_return=1
ASAN_CXXFLAGS:
default: -fsanitize=address
-fsanitize-address-use-after-scope
ASAN_LDFLAGS:
default: -fsanitize=address
# https://clang.llvm.org/docs/UndefinedBehaviorSanitizer.html
UBSAN_OPTIONS:
default: halt_on_error=1, print_stacktrace=1
UBSAN_CXXFLAGS:
default: -fsanitize=undefined
UBSAN_LDFLAGS:
default: -fsanitize=undefined
# https://clang.llvm.org/docs/MemorySanitizer.html
MSAN_OPTIONS:
default: "poison_in_dtor=1"
MSAN_CXXFLAGS:
default: -fsanitize=memory
-fsanitize-memory-track-origins
-fsanitize-memory-use-after-dtor
MSAN_LDFLAGS:
default: -fsanitize=memory
LSAN_DIR:
description: LSAN suppression directory
default: lsan
LSAN_FILE:
description: LSAN suppression file
default: lsan.supp
NO_FLAGS:
description: If true, do not set any CXXFLAGS or LDFLAGS.
default: false
runs:
using: 'composite'
steps:
- name: Env (Inputs)
run: |
echo DEBUG=${{inputs.DEBUG}} >> $GITHUB_ENV
echo REPOSITORY=${{inputs.REPOSITORY}} >> $GITHUB_ENV
echo BRANCH=${{inputs.BRANCH}} >> $GITHUB_ENV
echo CLANG_VER=${{inputs.CLANG_VER}} >> $GITHUB_ENV
echo LLVM_VER=${{inputs.LLVM_VER}} >> $GITHUB_ENV
echo HYPRE_VER=${{inputs.HYPRE_VER}} >> $GITHUB_ENV
echo METIS_VER=${{inputs.METIS_VER}} >> $GITHUB_ENV
echo CTEST=${{inputs.CTEST}} >> $GITHUB_ENV
echo ASAN_OPTIONS=${{inputs.ASAN_OPTIONS}} >> $GITHUB_ENV
echo UBSAN_OPTIONS=${{inputs.UBSAN_OPTIONS}} >> $GITHUB_ENV
echo MSAN_OPTIONS=${{inputs.MSAN_OPTIONS}} >> $GITHUB_ENV
echo LSAN_DIR=${{inputs.LSAN_DIR}} >> $GITHUB_ENV
echo LSAN_FILE=${{inputs.LSAN_FILE}} >> $GITHUB_ENV
echo ASAN_CXXFLAGS=${{inputs.ASAN_CXXFLAGS}} >> $GITHUB_ENV
echo ASAN_LDFLAGS=${{inputs.ASAN_LDFLAGS}} >> $GITHUB_ENV
echo UBSAN_CXXFLAGS=${{inputs.UBSAN_CXXFLAGS}} >> $GITHUB_ENV
echo UBSAN_LDFLAGS=${{inputs.UBSAN_LDFLAGS}} >> $GITHUB_ENV
echo MSAN_CXXFLAGS=${{inputs.MSAN_CXXFLAGS}} >> $GITHUB_ENV
echo MSAN_LDFLAGS=${{inputs.MSAN_LDFLAGS}} >> $GITHUB_ENV
shell: bash
- name: Env (dir)
run: |
echo LLVM_DIR=${{github.workspace}}/llvm >> $GITHUB_ENV
echo HYPRE_DIR=hypre-${{inputs.HYPRE_VER}} >> $GITHUB_ENV
echo METIS_DIR=metis-${{inputs.METIS_VER}} >> $GITHUB_ENV
shell: bash
- name: Env (bis)
run: |
echo CC=clang-${{inputs.CLANG_VER}} >> $GITHUB_ENV
echo CXX=clang++-${{inputs.CLANG_VER}} >> $GITHUB_ENV
echo LLVM_INC=${{env.LLVM_DIR}}/include/c++/v1 >> $GITHUB_ENV
echo LLVM_LIB=${{env.LLVM_DIR}}/lib >> $GITHUB_ENV
echo HYPRE_TGZ=v${{inputs.HYPRE_VER}}.tar.gz >> $GITHUB_ENV
echo METIS_TGZ=metis-${{inputs.METIS_VER}}.tar.gz >> $GITHUB_ENV
LSAN_SUPPRESSIONS="${{github.workspace}}/${{inputs.LSAN_DIR}}/${{inputs.LSAN_FILE}}"
echo "LSAN_OPTIONS=suppressions=$LSAN_SUPPRESSIONS" >> $GITHUB_ENV
shell: bash
- name: Env (ter)
if: ${{ inputs.NO_FLAGS != 'true' }}
run: |
echo LLVM_CXXFLAGS=-stdlib=libc++ -I${{env.LLVM_INC}} -Isystem${{env.LLVM_INC}} >> $GITHUB_ENV
echo LLVM_LDFLAGS=-L${{env.LLVM_LIB}} -lc++abi -Wl,-rpath,${{env.LLVM_LIB}} >> $GITHUB_ENV
shell: bash
- name: Env (quater)
if: ${{ inputs.NO_FLAGS != 'true' }}
run: |
echo CXXFLAGS=${{env.LLVM_CXXFLAGS}} >> $GITHUB_ENV
echo LDFLAGS=${{env.LLVM_LDFLAGS}} >> $GITHUB_ENV
shell: bash
-91
View File
@@ -1,91 +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.
---
name: 'MFEM Compilation'
description: 'MFEM Compilation'
inputs:
par:
description: 'Whether to build for parallel (true/false)'
default: false
sanitizer:
description: 'Sanitizer to use (asan, msan, ubsan)'
default: asan
runs:
using: 'composite'
steps:
- uses: ./.github/actions/sanitize/config
- uses: actions/cache@v4
if: ${{env.DEBUG == 'true'}}
id: debug
with:
path: mfem/build
key: build-${{inputs.par}}-${{inputs.sanitizer}}
- uses: ./.github/actions/sanitize/setup
if: ${{steps.debug.outputs.cache-hit != 'true'}}
with:
par: ${{inputs.par}}
sanitizer: ${{inputs.sanitizer}}
- name: Build with ASAN
if: inputs.sanitizer == 'asan'
run: echo CXXFLAGS=${{env.CXXFLAGS}} ${{env.ASAN_CXXFLAGS}} >> $GITHUB_ENV
shell: bash
- name: Build with MSAN
if: inputs.sanitizer == 'msan'
run: echo CXXFLAGS=${{env.CXXFLAGS}} ${{env.MSAN_CXXFLAGS}} >> $GITHUB_ENV
shell: bash
- name: Build with UBSAN
if: inputs.sanitizer == 'ubsan'
run: echo CXXFLAGS=${{env.CXXFLAGS}} ${{env.UBSAN_CXXFLAGS}} >> $GITHUB_ENV
shell: bash
- uses: mfem/github-actions/build-mfem@v2.5
if: ${{steps.debug.outputs.cache-hit != 'true'}}
env:
CXXFLAGS: ${{env.CXXFLAGS}}
LDFLAGS: ${{env.LDFLAGS}}
with:
mpi: ${{inputs.par == 'false' && 'seq' || 'par'}}
mfem-dir: mfem
os: ${{runner.os}}
library-only: true
build-system: cmake
hypre-dir: ${{env.HYPRE_DIR}}
metis-dir: ${{env.METIS_DIR}}
config-options: >-
-GNinja
-DMPICXX=${{env.CXX}}
-DCMAKE_CXX_STANDARD=17
-DMFEM_USE_MEMALLOC=OFF
-DCMAKE_BUILD_TYPE=Release
-DCMAKE_VERBOSE_MAKEFILE=ON
-DCMAKE_CXX_COMPILER=${{env.CXX}}
-DCMAKE_CXX_FLAGS_RELEASE='-g -O1 -fno-omit-frame-pointer'
- name: Delete object files
if: ${{steps.debug.outputs.cache-hit != 'true'}}
working-directory: mfem/build
run: find . -type f -name '*.o' -delete
shell: bash
- uses: actions/upload-artifact@v4
with:
name: build-${{inputs.par}}-${{inputs.sanitizer}}
path: mfem/build
if-no-files-found: error
retention-days: 1
overwrite: false
-33
View File
@@ -1,33 +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.
---
name: 'Install MPI'
description: 'Installs MPI and set up its environment variables'
runs:
using: 'composite'
steps:
- name: Install
run: sudo apt-get install openmpi-bin libopenmpi-dev
shell: bash
- name: Env
run: |
echo PRTE_MCA_rmaps_default_mapping_policy=:oversubscribe >> $GITHUB_ENV
echo MPI_INC=$(mpicxx --showme:compile) >> $GITHUB_ENV
echo MPI_LIB=$(mpicxx --showme:link) >> $GITHUB_ENV
shell: bash
- name: Env (bis)
run: |
echo CXXFLAGS=${{env.CXXFLAGS}} ${{env.MPI_INC}} >> $GITHUB_ENV
echo LDFLAGS=${{env.LDFLAGS}} ${{env.MPI_LIB}} >> $GITHUB_ENV
shell: bash
@@ -1,71 +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.
---
name: 'Restore state'
description: 'Restore state to be able to run checks, tests'
inputs:
par:
description: 'Whether to build for parallel (true/false)'
default: false
sanitizer:
description: 'Sanitizer to use (asan, msan, ubsan)'
default: asan
cache-path:
description: 'path to what needs to be restored'
default: none
cache-skip:
description: 'Skip cache restoration'
default: false
outputs:
cache-hit:
description: 'Output from a specific step'
value: ${{steps.debug.outputs.cache-hit}}
runs:
using: 'composite'
steps:
- uses: ./.github/actions/sanitize/config
- uses: actions/cache@v4
if: ${{env.DEBUG == 'true' && inputs.cache-skip != 'true'}}
id: debug
with:
path: ${{inputs.cache-path}}
key: ${{github.job}}-${{inputs.par}}-${{inputs.sanitizer}}
- uses: ./.github/actions/sanitize/setup
if: ${{steps.debug.outputs.cache-hit != 'true'}}
with:
par: ${{inputs.par}}
sanitizer: ${{inputs.sanitizer}}
- uses: actions/download-artifact@v4
with:
name: build-${{inputs.par}}-${{inputs.sanitizer}}
path: mfem/build
- name: Ninja Patch
working-directory: mfem/build
run: |
sed -i -e 's/CXX_STATIC_LIBRARY_LINKER__mfem_Release.*/CUSTOM_COMMAND/' build.ninja
sed -i -e '/build tests\/unit\/all:/ s/tests\/unit\/[^ ]*unit_tests[^ ]*//g' build.ninja
sed -i -e '/^add_test(\[=\[\(unit_tests\|punit_tests\)\]=\]/ s/)/ "--input-file .\/list-test-names-${{matrix.tag}}" "--min-duration 1")/' tests/unit/CTestTestfile.cmake
shell: bash
- name: Copy Data
if: ${{steps.debug.outputs.cache-hit != 'true'}}
working-directory: mfem/build
run: |
ninja cmake_object_order_depends_target_unit_tests
cp -pR ../tests/unit/data tests/unit
shell: bash
-64
View File
@@ -1,64 +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.
---
name: 'Setup state'
description: 'Sets up the state to be able to run build & run'
inputs:
par:
description: 'Whether to build for parallel (true/false)'
default: false
sanitizer:
description: 'Sanitizer to use (asan, msan, ubsan)'
default: asan
runs:
using: 'composite'
steps:
- uses: actions/cache/restore@v4 # Cache for LLVM libcxx
with:
path: ${{env.LLVM_DIR}}
fail-on-cache-miss: true
key: build-libcxx-${{env.LLVM_VER}}-${{inputs.sanitizer}}
- uses: ./.github/actions/sanitize/mpi
if: ${{inputs.par == 'true'}}
- uses: actions/cache/restore@v4 # Cache for Hypre
if: ${{inputs.par == 'true'}}
with:
path: ${{env.HYPRE_DIR}}
fail-on-cache-miss: true
key: ${{runner.os}}-ompi-build-${{env.HYPRE_DIR}}-int32-fp64-v2.5
- uses: actions/cache/restore@v4 # Cache for Metis
if: ${{inputs.par == 'true'}}
with:
path: ${{env.METIS_DIR}}
fail-on-cache-miss: true
key: ${{runner.os}}-build-${{env.METIS_DIR}}-v2.5
- name: Hypre/Metis links
if: ${{inputs.par == 'true'}}
run: ln -s -f ${{env.HYPRE_DIR}} hypre && ln -s -f ${{env.METIS_DIR}} metis-4.0
shell: bash
- uses: actions/cache/restore@v4 # Cache for LSAN suppression file
with:
path: ${{env.LSAN_DIR}}
fail-on-cache-miss: true
key: build-lsan-suppression-file
- uses: actions/checkout@v4 # Checkout the repository
with:
path: mfem
# ref: ${{env.BRANCH}}
# repository: ${{env.REPOSITORY}}
+7 -26
View File
@@ -7,17 +7,18 @@
https://mfem.org
This directory contains the GitHub CI scripts for MFEM.
Note that some of these scripts use the shared MFEM GitHub Actions from the external mfem/github-actions repository:
<https://github.com/mfem/github-actions>
https://github.com/mfem/github-actions
For a particular action, e.g. `mfem/github-actions/build-mfem@v2.5`, the `v2.5` suffix denotes the branch in the above from which the action is taken.
For a particular action, e.g. `mfem/github-actions/build-mfem@v2.1`, the `v2.1` suffix denotes the branch in the above from which the action is taken.
The current CI workflows are:
## `repo-check.yml`
### `repo-check.yml`
Runs a number of static repository-level sanity checks.
@@ -29,39 +30,19 @@ Runs a number of static repository-level sanity checks.
- `branch-history` guards against accidental commits of large files using the `--history` option of the `config/githooks/pre-push` script.
## `mfem-analysis.yml` (`build-analysis`)
### `mfem-analysis.yml` (`build-analysis`)
Checks if the code builds and satisfies minimal requirements.
- `gitignore` builds hypre, METIS, and MFEM using `mfem/github-actions/build-hypre`, `mfem/github-actions/build-metis`, and `mfem/github-actions/build-mfem` and checks for correct `.gitignore` settings by running the `tests/scripts/gitignore` script.
## `builds-and-tests.yml`
### `builds-and-tests.yml`
Runs a matrix of builds and tests runs with different compilers, OS, mfem/hypre settings, etc. Also processes and upload Codecov reports.
Uses the following GitHub Actions from <https://github.com/mfem/github-actions>:
Uses the following GitHub Actions from https://github.com/mfem/github-actions:
- `mfem/github-actions/build-hypre`
- `mfem/github-actions/build-metis`
- `mfem/github-actions/build-mfem`
- `mfem/github-actions/upload-coverage`
## Sanitizer Workflow for MFEM Verification
This workflow validates MFEM unit tests, examples, and miniapps using sanitizer tools.
- `sanitizers.yml` orchestrates:
- Building and caching dependencies: HYPRE, METIS, LSAN suppression file, and LLVM libcxx.
- Launching fine-grained jobs for serial (ASAN, MSAN, UBSAN) and parallel (ASAN, UBSAN) sanitizers.
- `sanitize-tests.yml` is a reusable workflow accepting `par` mode (`true` for parallel) and `sanitizer` (ASAN, MSAN, or UBSAN) as inputs. It executes the following jobs:
- **Build**: Compiles the MFEM library with specified parallel and sanitizer settings.
- **Check**: Runs verification checks.
- Parallel jobs to test the following: **Examples**, **Miniapps** and **Unit tests**
The workflow leverages composite actions in `.github/actions/sanitize/`:
- `config`: Centralizes settings for the sanitizer workflow.
- `mfem`: Manages the MFEM library build process.
- `mpi`: Installs MPI and applies additional compilation flags.
- `restore`: Restores the testing environment state.
- `setup`: Builds or restores cached dependencies.
+4 -4
View File
@@ -289,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
+69
View File
@@ -0,0 +1,69 @@
# 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.
name: "Sanitizer"
permissions:
actions: write
on:
push:
branches:
- master
- next
pull_request:
workflow_dispatch:
concurrency:
group: ${{ github.workflow }}-${{ github.ref }}
cancel-in-progress: true
jobs:
Serial:
runs-on: ubuntu-24.04
steps:
- name: MFEM Checkout
uses: actions/checkout@v4
with:
path: mfem
- name: MFEM Build
uses: mfem/github-actions/build-mfem@v2.5
with:
os: ${{ runner.os }}
target: opt
mpi: seq
hypre-dir: unused-hypre-dir
metis-dir: unused-metis-dir
mfem-dir: mfem
build-system: make
library-only: false
config-options:
CXX="clang++-18"
CXXFLAGS="-g -O1 -std=c++17
-fsanitize=address
-fno-omit-frame-pointer
-fsanitize-address-use-after-scope"
- name: MFEM Info
working-directory: mfem
run: make info
- name: MFEM Sanitize
working-directory: mfem
run:
ASAN_OPTIONS="detect_leaks=1,
strict_init_order=1,
strict_string_checks=1,
check_initialization_order=1,
detect_stack_use_after_return=1"
make test
@@ -1,39 +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.
---
name: build-hypre
on:
workflow_call:
jobs:
build-hypre:
runs-on: ubuntu-latest
name: 2.19.0
steps:
- uses: actions/checkout@v4
- uses: ./.github/actions/sanitize/config
- name: Cache
id: cache
uses: actions/cache@v4
with:
path: ${{env.HYPRE_DIR}}
key: ${{runner.os}}-ompi-build-${{env.HYPRE_DIR}}-int32-fp64-v2.5
- name: Setup
if: steps.cache.outputs.cache-hit != 'true'
uses: ./.github/actions/sanitize/mpi
- name: Build
if: steps.cache.outputs.cache-hit != 'true'
uses: mfem/github-actions/build-hypre@v2.5
with:
archive: ${{env.HYPRE_TGZ}}
dir: ${{env.HYPRE_DIR}}
target: int32
precision: fp64
build-system: make
@@ -1,76 +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.
---
name: build-libcxx
on:
workflow_call:
jobs:
build-llvm-libcxx:
runs-on: ubuntu-latest
strategy:
matrix:
sanitizer: [asan, msan, ubsan]
include:
- sanitizer: asan
llvm_use_sanitizer: "Address"
- sanitizer: msan
llvm_use_sanitizer: "MemoryWithOrigins"
- sanitizer: ubsan
llvm_use_sanitizer: "Undefined"
name: ${{matrix.sanitizer}}
steps:
- uses: actions/checkout@v4
- uses: ./.github/actions/sanitize/config
with:
NO_FLAGS: true
- name: Cache
id: cache
uses: actions/cache@v4
with:
path: ${{env.LLVM_DIR}}
key: build-libcxx-${{env.LLVM_VER}}-${{matrix.sanitizer}}
- name: Clone
if: ${{ steps.cache.outputs.cache-hit != 'true' }}
run: >
git clone --filter=blob:none --depth=1
--branch llvmorg-${{env.LLVM_VER}}
--no-checkout https://github.com/llvm/llvm-project.git llvm-project
- name: Checkout
if: ${{ steps.cache.outputs.cache-hit != 'true' }}
working-directory: llvm-project
run: |
git sparse-checkout set --cone
git checkout llvmorg-${{env.LLVM_VER}}
git sparse-checkout set cmake llvm/cmake runtimes libcxx libcxxabi
- name: Mkdir
if: ${{ steps.cache.outputs.cache-hit != 'true' }}
run: mkdir ${{env.LLVM_DIR}}
- name: CMake
if: ${{ steps.cache.outputs.cache-hit != 'true' }}
working-directory: ${{env.LLVM_DIR}}
run: >
VERBOSE=1
cmake -GNinja ../llvm-project/runtimes/
-DCMAKE_C_COMPILER=${{env.CC}}
-DCMAKE_CXX_COMPILER=${{env.CXX}}
-DCMAKE_BUILD_TYPE=RelWithDebInfo
-DCMAKE_INSTALL_PREFIX=/usr
-DLLVM_USE_SANITIZER=${{matrix.llvm_use_sanitizer}}
-DLLVM_BUILD_32_BITS=OFF
-DLIBCXXABI_USE_LLVM_UNWINDER=OFF
-DLLVM_INCLUDE_TESTS=OFF
-DLIBCXX_INCLUDE_TESTS=OFF
-DLIBCXX_INCLUDE_BENCHMARKS=OFF
-DLLVM_ENABLE_RUNTIMES='libcxx;libcxxabi'
- name: Build
if: ${{ steps.cache.outputs.cache-hit != 'true' }}
working-directory: ${{env.LLVM_DIR}}
run: cmake --build . -- cxx cxxabi
-38
View File
@@ -1,38 +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.
---
name: build-file-lsan
on:
workflow_call:
jobs:
build-file-lsan:
runs-on: ubuntu-latest
name: lsan.supp
steps:
- uses: actions/checkout@v4
- uses: ./.github/actions/sanitize/config
- name: Cache
id: cache
uses: actions/cache@v4
with:
path: ${{env.LSAN_DIR}}
key: build-lsan-suppression-file
- name: Setup
if: steps.cache.outputs.cache-hit != 'true'
run: |
mkdir -p ${{env.LSAN_DIR}}
cat << EOF > ${{env.LSAN_DIR}}/${{env.LSAN_FILE}}
leak:libevent_core-2.1.so
leak:ompi_mpi_finalize
leak:ompi_mpi_init
leak:PMPI_Init
leak:strdup
EOF
@@ -1,36 +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.
---
name: build-metis
on:
workflow_call:
jobs:
build-metis:
runs-on: ubuntu-latest
name: 4.0.3
steps:
- uses: actions/checkout@v4
- uses: ./.github/actions/sanitize/config
- name: Cache
id: cache
uses: actions/cache@v4
with:
path: ${{env.METIS_DIR}}
key: ${{runner.os}}-build-${{env.METIS_DIR}}-v2.5
- name: Setup
if: steps.cache.outputs.cache-hit != 'true'
uses: ./.github/actions/sanitize/mpi
- name: Build
if: steps.cache.outputs.cache-hit != 'true'
uses: mfem/github-actions/build-metis@v2.5
with:
archive: ${{env.METIS_TGZ}}
dir: ${{env.METIS_DIR}}
-197
View File
@@ -1,197 +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.
---
name: Sanitize
on:
workflow_call:
inputs:
par:
description: 'Whether to build for parallel (true/false)'
required: false
default: false
type: boolean
sanitizer:
description: 'Sanitizer to use (asan, msan, ubsan)'
required: true
default: asan
type: string
jobs:
build:
runs-on: ubuntu-latest
steps:
- uses: actions/checkout@v4
- uses: ./.github/actions/sanitize/mfem
with:
par: ${{inputs.par}}
sanitizer: ${{inputs.sanitizer}}
check:
needs: [build]
runs-on: ubuntu-latest
env:
ex: ${{inputs.par && 'ex1p' || 'ex1'}}
steps:
- uses: actions/checkout@v4
- uses: ./.github/actions/sanitize/restore
id: restore
with:
par: ${{inputs.par}}
sanitizer: ${{inputs.sanitizer}}
cache-path: mfem/build/examples/${{env.ex}}
- name: MFEM Check
if: ${{steps.restore.outputs.cache-hit != 'true'}}
working-directory: mfem/build
run: ninja -v check
examples:
needs: [check]
runs-on: ubuntu-latest
env:
exclude: ${{inputs.par && '-E "_ser"' || ''}}
steps:
- uses: actions/checkout@v4
- uses: ./.github/actions/sanitize/restore
id: restore
with:
par: ${{inputs.par}}
sanitizer: ${{inputs.sanitizer}}
cache-path: mfem/build/examples/ex1
- name: Build Examples
if: ${{steps.restore.outputs.cache-hit != 'true'}}
working-directory: mfem/build
run: ninja -v examples
- name: Test Examples
if: ${{steps.restore.outputs.cache-hit != 'true'}}
working-directory: mfem/build
run: |
${{env.CTEST}} examples ${{env.exclude}} --show-only
${{env.CTEST}} examples ${{env.exclude}}
miniapps:
needs: [check]
runs-on: ubuntu-latest
env:
exclude: ${{inputs.par && '-E "_ser"' || ''}}
steps:
- uses: actions/checkout@v4
- uses: ./.github/actions/sanitize/restore
id: restore
with:
par: ${{inputs.par}}
sanitizer: ${{inputs.sanitizer}}
cache-path: mfem/build/miniapps/meshing/minimal-surface
- name: Build Miniapps
if: ${{steps.restore.outputs.cache-hit != 'true'}}
working-directory: mfem/build
run: ninja -v miniapps
- name: Test Miniapps
if: ${{steps.restore.outputs.cache-hit != 'true'}}
working-directory: mfem/build
run: |
${{env.CTEST}} miniapps ${{env.exclude}} --show-only
${{env.CTEST}} miniapps ${{env.exclude}}
tests-miniapps:
needs: [check]
runs-on: ubuntu-latest
env:
run: ${{inputs.par && '-R "_cpu_np"' || ''}}
exclude: ${{inputs.par && '"unit_tests|debug"' || '"^unit_tests$|debug"'}}
steps:
- uses: actions/checkout@v4
- uses: ./.github/actions/sanitize/restore
id: restore
with:
par: ${{inputs.par}}
sanitizer: ${{inputs.sanitizer}}
cache-path: mfem/build/tests/unit/sedov_tests_cpu
- name: Build Tests Unit Miniapps
if: ${{steps.restore.outputs.cache-hit != 'true'}}
working-directory: mfem/build
run: ninja -v tests/unit/all
- name: Run Tests Unit Miniapps
if: ${{steps.restore.outputs.cache-hit != 'true'}}
working-directory: mfem/build
run: |
${{env.CTEST}} tests/unit -E ${{env.exclude}} ${{env.run}} --show-only
${{env.CTEST}} tests/unit -E ${{env.exclude}} ${{env.run}}
tests-unit-build:
needs: [check]
runs-on: ubuntu-latest
env:
unit_tests: ${{inputs.par && 'punit_tests' || 'unit_tests'}}
steps:
- uses: actions/checkout@v4
- uses: ./.github/actions/sanitize/restore
id: restore
with:
par: ${{inputs.par}}
sanitizer: ${{inputs.sanitizer}}
cache-path: mfem/build/tests/unit/${{env.unit_tests}}
- name: Build Unit Tests
if: ${{steps.restore.outputs.cache-hit != 'true'}}
working-directory: mfem/build
run: ninja -v ${{env.unit_tests}}
- name: Delete object files
if: ${{steps.restore.outputs.cache-hit != 'true'}}
working-directory: mfem/build/tests/unit
run: find . -type f -name '*.o' -delete
- uses: actions/upload-artifact@v4
with:
name: tests-${{inputs.par}}-${{inputs.sanitizer}}
path: mfem/build/tests/unit/${{env.unit_tests}}
if-no-files-found: error
retention-days: 1
overwrite: false
tests-unit-run:
needs: [tests-unit-build]
runs-on: ubuntu-latest
strategy:
matrix:
tag: [0, 1, 2, 3]
name: tests-unit-run-${{matrix.tag}}
env:
unit_tests: ${{inputs.par && 'punit_tests' || 'unit_tests'}}
np: ${{inputs.par && '_np=2' || ''}}
steps:
- uses: actions/checkout@v4
- uses: ./.github/actions/sanitize/restore
id: restore
with:
par: ${{inputs.par}}
sanitizer: ${{inputs.sanitizer}}
cache-path: mfem/build/tests/unit/${{env.unit_tests}}
- uses: actions/download-artifact@v4
if: ${{steps.restore.outputs.cache-hit != 'true'}}
with:
name: tests-${{inputs.par}}-${{inputs.sanitizer}}
path: mfem/build/tests/unit
- name: Split Unit Tests
if: ${{steps.restore.outputs.cache-hit != 'true'}}
working-directory: mfem/build/tests/unit
run: |
chmod 755 ${{env.unit_tests}}
./${{env.unit_tests}} --list-test-names-only | tail -n +2 > list-test-names
shuf list-test-names -o list-test-names
split --verbose -n l/4 -d -a 1 list-test-names list-test-names-
- name: Cat Unit Tests ${{matrix.tag}}
if: ${{steps.restore.outputs.cache-hit != 'true'}}
working-directory: mfem/build/tests/unit
run: cat list-test-names-${{matrix.tag}}
- name: Run Unit Tests ${{matrix.tag}}
if: ${{steps.restore.outputs.cache-hit != 'true'}}
working-directory: mfem/build
run: |
${{env.CTEST}} tests/unit -R "${{env.unit_tests}}${{env.np}}" --show-only
${{env.CTEST}} tests/unit -R "${{env.unit_tests}}${{env.np}}"
-73
View File
@@ -1,73 +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.
---
name: Sanitizers
permissions:
actions: write
on:
push:
branches: ["master", "next"]
pull_request:
workflow_dispatch:
concurrency:
group: ${{github.workflow}}-${{github.ref}}
cancel-in-progress: true
jobs:
# Build steps for dependencies
build-hypre:
uses: ./.github/workflows/sanitize-build-hypre.yml
build-metis:
uses: ./.github/workflows/sanitize-build-metis.yml
build-lsan:
uses: ./.github/workflows/sanitize-build-lsan.yml
build-libcxx:
uses: ./.github/workflows/sanitize-build-libcxx.yml
# Serial sanitizers: asan, msan, ubsan
seq-asan:
needs: [build-libcxx]
uses: ./.github/workflows/sanitize-tests.yml
with:
sanitizer: asan
seq-msan:
needs: [build-libcxx]
uses: ./.github/workflows/sanitize-tests.yml
with:
sanitizer: msan
seq-ubsan:
needs: [build-libcxx]
uses: ./.github/workflows/sanitize-tests.yml
with:
sanitizer: ubsan
# Parallel sanitizers: asan, ubsan
par-asan:
needs: [build-libcxx, build-hypre, build-metis]
uses: ./.github/workflows/sanitize-tests.yml
with:
par: true
sanitizer: asan
par-ubsan:
needs: [build-libcxx, build-hypre, build-metis]
uses: ./.github/workflows/sanitize-tests.yml
with:
par: true
sanitizer: ubsan
+9 -10
View File
@@ -19,6 +19,9 @@ CMakeFiles/
# Clangd server cache
*.cache*
# VSCode configuration
/.vscode/
# Backup files
*~
@@ -211,7 +214,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 +270,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 +280,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 +303,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 +321,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,14 +412,12 @@ 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
tests/unit/punit_tests
tests/unit/gpu_unit_tests
tests/unit/pgpu_unit_tests
tests/unit/cunit_tests
tests/unit/pcunit_tests
tests/unit/sedov_tests_*
tests/unit/psedov_tests_*
tests/unit/tmop_pa_tests_*
+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
+1 -33
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
@@ -44,13 +39,6 @@ GPU computing
set. This is most often used for setting constant essential boundary
conditions. A new function Vector::SetSubVectorHost has been added in cases
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
-------------------------------------
@@ -60,33 +48,13 @@ New and updated examples and miniapps
operators as smoothers.
These miniapps can be found in `miniapps/diag-smoothers`.
API changes
API changes:
-----------
- mfem::internal::tensor and mfem::internal::dual have been moved to
mfem::future::tensor and mfem::future::dual.
- API addition: in class `Operator`, added virtual functions: `AbsMult`, and
`AbsMultTranspose`; in class `Vector`, added `Abs` and `Pow`.
Miscellaneous
-------------
- Added the "gpu", "raja-gpu", and "ceed-gpu" backend aliases/shortcuts which
automatically select between CUDA or HIP.
- The CUDA-specific names used by some of the unit tests like 'cunit_tests' and
'pcunit_tests' were replaced by names using 'gpu' instead of 'c' (short for
CUDA) or 'cuda'. These tests automatically run the CUDA/HIP tests based on the
MFEM build configuration.
- Added the option to enable GPU-aware MPI in MFEM using the environment
variable 'MFEM_GPU_AWARE_MPI' set to any value. Setting this environment
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++)
+1 -3
View File
@@ -78,7 +78,6 @@ 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));
@@ -101,8 +100,7 @@ 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,7 +88,6 @@ 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));
-3
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
@@ -249,7 +247,6 @@ set(HDRS
nonlinearform_ext.hpp
nonlininteg.hpp
qfunction.hpp
qinterp/det.hpp
qinterp/eval.hpp
qinterp/eval_hdiv.hpp
qinterp/grad.hpp
-1
View File
@@ -515,7 +515,6 @@ struct InvTNewtonSolver<Geometry::SEGMENT, SDim, SType, max_team_x>
phys_tol += pptr[idx + d * npts] * pptr[idx + d * npts];
}
phys_tol = fmax(phys_rtol * phys_rtol, phys_tol * phys_rtol * phys_rtol);
hit_bdr[0] = prev_hit_bdr[0] = false;
}
// for each iteration
while (true)
+64 -34
View File
@@ -255,8 +255,6 @@ 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);
@@ -268,7 +266,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
@@ -327,14 +329,50 @@ 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();
@@ -391,7 +429,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)
{
@@ -412,7 +450,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);
@@ -438,7 +476,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);
}
}
@@ -451,7 +489,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);
}
@@ -491,8 +529,6 @@ 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();
@@ -552,7 +588,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)
{
@@ -654,8 +690,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)
{
@@ -663,14 +699,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);
@@ -693,7 +727,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);
@@ -740,14 +774,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);
}
@@ -771,7 +804,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);
@@ -865,7 +898,6 @@ EABilinearFormExtension::EABilinearFormExtension(BilinearForm *form)
void EABilinearFormExtension::Assemble()
{
MFEM_PERF_FUNCTION;
SetupRestrictionOperators(L2FaceValues::SingleValued);
ne = trial_fes->GetMesh()->GetNE();
@@ -888,8 +920,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;
@@ -922,7 +953,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);
}
}
}
@@ -951,7 +982,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);
}
}
}
@@ -1000,7 +1031,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);
}
}
}
@@ -1414,7 +1445,6 @@ FABilinearFormExtension::FABilinearFormExtension(BilinearForm *form)
void FABilinearFormExtension::Assemble()
{
MFEM_PERF_FUNCTION;
EABilinearFormExtension::Assemble();
FiniteElementSpace &fes = *a->FESpace();
int width = fes.GetVSize();
+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;
+10 -10
View File
@@ -812,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 );
}
@@ -884,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 );
}
@@ -919,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 );
}
@@ -1600,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 );
}
@@ -1635,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 );
}
@@ -1669,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 );
}
@@ -1760,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 );
}
@@ -1793,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 );
}
@@ -1832,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
);
@@ -1973,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 );
}
-1
View File
@@ -50,7 +50,6 @@ ElementTransformation *RefinedToCoarse(
void Coefficient::Project(QuadratureFunction &qf)
{
MFEM_PERF_FUNCTION;
QuadratureSpaceBase &qspace = *qf.GetSpace();
const int ne = qspace.GetNE();
Vector values;
+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;
}
}
+10 -12
View File
@@ -764,9 +764,9 @@ ParaViewDataCollectionBase::ParaViewDataCollectionBase(
{
cycle = 0;
#ifdef MFEM_USE_ZLIB
// If we have zlib, enable compression. Otherwise, compression is disabled in
// the DataCollection base class constructor.
compression = true;
compression = true; // if we have zlib, enable compression
#else
compression = false; // otherwise, disable compression
#endif
}
@@ -784,8 +784,13 @@ void ParaViewDataCollectionBase::SetCompressionLevel(int compression_level_)
{
MFEM_ASSERT(compression_level_ >= -1 && compression_level_ <= 9,
"Compression level must be between -1 and 9 (inclusive).");
if (compression_level_ != 0) { SetCompression(true);}
compression_level = compression_level_;
compression = compression_level_ != 0;
}
void ParaViewDataCollectionBase::SetCompression(bool compression_)
{
compression = compression_;
}
int ParaViewDataCollectionBase::GetCompressionLevel() const
@@ -1169,14 +1174,7 @@ const char *ParaViewDataCollection::GetDataTypeString() const
ParaViewHDFDataCollection::ParaViewHDFDataCollection(
const std::string &collection_name, Mesh *mesh)
: ParaViewDataCollectionBase(collection_name, mesh)
{
compression = true;
}
void ParaViewHDFDataCollection::SetCompression(bool compression_)
{
compression = compression_;
}
{ }
void ParaViewHDFDataCollection::EnsureVTKHDF()
{
+7 -6
View File
@@ -537,6 +537,13 @@ public:
/// Any nonzero compression level will enable compression.
void SetCompressionLevel(int compression_level_);
/// @brief Enable or disable zlib compression.
///
/// If the input is true, use the default zlib compression level (unless the
/// compression level has previously been set by calling
/// SetCompressionLevel()).
void SetCompression(bool compression_) override;
/// @brief Sets whether or not to output the data as high-order elements
/// (false by default).
///
@@ -626,12 +633,6 @@ public:
ParaViewHDFDataCollection(const std::string& collection_name,
Mesh *mesh_ = nullptr);
/// @brief Enable or disable compression.
///
/// The compression level can be set with SetCompressionLevel()). VTKHDF
/// compression does not require MFEM to be compiled with zlib support.
void SetCompression(bool compression_) override;
/// Save the collection.
void Save() override;
-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);
+8 -9
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());
}
@@ -1708,7 +1708,6 @@ std::array<DofToQuadMap, N> load_dtq_mem(
const auto B = Reshape(&dtq[i].B[0], nqp_b, dim_b, ndof_b);
auto mem_Bi = Reshape(reinterpret_cast<real_t *>(mem) + offset, nqp_b, dim_b,
ndof_b);
MFEM_FOREACH_THREAD(q, x, nqp_b)
{
MFEM_FOREACH_THREAD(d, y, ndof_b)
@@ -2159,7 +2158,7 @@ template <
std::size_t... Is>
std::array<DofToQuadMap, N> create_dtq_maps_impl(
field_operator_ts &fops,
std::vector<const DofToQuad*> &dtqs,
std::vector<const DofToQuad*> dtqs,
const std::array<int, N> &field_map,
std::index_sequence<Is...>)
{
@@ -2244,7 +2243,7 @@ template <
std::size_t num_fields>
std::array<DofToQuadMap, num_fields> create_dtq_maps(
field_operator_ts &fops,
std::vector<const DofToQuad*> &dtqmaps,
std::vector<const DofToQuad*> dtqmaps,
const std::array<int, num_fields> &to_field_map)
{
return create_dtq_maps_impl<entity_t>(
-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
+7 -21
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),
@@ -101,10 +98,7 @@ FiniteElementSpace::FiniteElementSpace(const FiniteElementSpace &orig,
FiniteElementSpace::FiniteElementSpace(Mesh *mesh,
const FiniteElementCollection *fec,
int vdim, int ordering)
{
MFEM_PERF_FUNCTION;
Constructor(mesh, NULL, fec, vdim, ordering);
}
{ Constructor(mesh, NULL, fec, vdim, ordering); }
FiniteElementSpace::FiniteElementSpace(Mesh *mesh, NURBSExtension *ext,
const FiniteElementCollection *fec,
@@ -396,8 +390,6 @@ void FiniteElementSpace::BuildElementToDofTable() const
{
if (elem_dof) { return; }
MFEM_PERF_FUNCTION;
// TODO: can we call GetElementDofs only once per element?
Table *el_dof = new Table;
Table *el_fos = (mesh->Dimension() > 2) ? (new Table) : NULL;
@@ -2753,8 +2745,6 @@ void FiniteElementSpace::BuildNURBSFaceToDofTable() const
void FiniteElementSpace::Construct()
{
MFEM_PERF_FUNCTION;
// This method should be used only for non-NURBS spaces.
MFEM_VERIFY(!NURBSext, "internal error");
@@ -4254,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).
+2 -121
View File
@@ -19,7 +19,6 @@
#include "../mesh/nurbs.hpp"
#include "../mesh/vtkhdf.hpp"
#include "../general/text.hpp"
#include "../general/reducers.hpp"
#ifdef MFEM_USE_MPI
#include "pfespace.hpp"
@@ -3327,126 +3326,8 @@ real_t GridFunction::ComputeLpError(const real_t p, Coefficient &exsol,
const IntegrationRule *irs[],
const Array<int> *elems) const
{
MFEM_PERF_FUNCTION;
MFEM_VERIFY(fes->GetVDim() == 1, "invalid vector dimension!");
real_t error = 0.0;
bool device_eval = true;
// TODO: check for cases that are not supported on device:
// * mixed meshes
// * meshes with non-tensor-product elements can have negative weights
// * variable orders
// * weight is not NULL
// * elems is not NULL
// * map type is not VALUE
// * ...
Mesh *mesh = fes->GetMesh();
const FiniteElement *fe = fes->GetTypicalFE();
if (mesh->GetNumGeometries(mesh->Dimension()) > 1 ||
(mesh->Dimension() > 1 && mesh->MeshGenerator() != 2) ||
fes->IsVariableOrder() ||
weight != nullptr ||
elems != nullptr ||
fe->GetMapType() != FiniteElement::MapType::VALUE)
{
device_eval = false;
}
if (device_eval)
{
Geometry::Type geom = mesh->GetTypicalElementGeometry();
const IntegrationRule *ir_p;
if (irs)
{
ir_p = irs[geom];
}
else
{
int intorder = 2*fe->GetOrder() + 3; // <----------
ir_p = &(IntRules.Get(geom, intorder));
}
const IntegrationRule &ir = *ir_p;
QuadratureSpace qs(*mesh, ir);
CoefficientVector coeff(exsol, qs, CoefficientStorage::FULL);
const QVectorLayout ql = QVectorLayout::byNODES;
const MemoryType d_mt = MemoryType::DEFAULT;
Vector q_vals;
// TODO: make this a method
{
// const FiniteElement *fe = fes->GetTypicalFE();
const int vdim = fes->GetVDim();
const int NE = fes->GetNE();
const int ND = fe->GetDof();
const int NQ = ir.GetNPoints();
MemoryType my_d_mt = (d_mt != MemoryType::DEFAULT) ? d_mt :
Device::GetDeviceMemoryType();
// byNODES : NQPT x VDIM x NE
// byVDIM : VDIM x NQPT x NE
q_vals.SetSize(vdim*NQ*NE, my_d_mt);
const QuadratureInterpolator &qi = *fes->GetQuadratureInterpolator(ir);
qi.SetOutputLayout(ql);
const bool use_tensor_products = UsesTensorBasis(*fes);
qi.DisableTensorProducts(!use_tensor_products);
const ElementDofOrdering e_ordering =
use_tensor_products ?
ElementDofOrdering::LEXICOGRAPHIC :
ElementDofOrdering::NATIVE;
const Operator *elem_restr = fes->GetElementRestriction(e_ordering);
if (fe->GetMapType() == FiniteElement::MapType::INTEGRAL)
{
// Pre-compute the geometric factors in order to set the desired
// MemoryType they use:
fes->GetMesh()->GetGeometricFactors(
ir, GeometricFactors::DETERMINANTS, my_d_mt);
}
if (elem_restr)
{
Vector f_e(vdim*ND*NE, my_d_mt);
elem_restr->Mult(*this, f_e);
qi.PhysValues(f_e, q_vals);
}
else
{
qi.PhysValues(*this, q_vals);
}
}
const real_t *exact_d = coeff.Read();
const real_t *gridf_d = q_vals.Read();
// FIXME: reuse the workspace vector from vector.cpp?
static Array<real_t> workspace;
if (p < infinity())
{
MemoryType my_d_mt = (d_mt != MemoryType::DEFAULT) ? d_mt :
Device::GetDeviceMemoryType();
const GeometricFactors *geom_factors =
fes->GetMesh()->GetGeometricFactors(
ir, GeometricFactors::DETERMINANTS, my_d_mt);
const real_t *detJ_d = geom_factors->detJ.Read();
const real_t *w_d = ir.GetWeights().Read();
const int NQ = ir.GetNPoints();
mfem::reduce(q_vals.Size(), error,
[=] MFEM_HOST_DEVICE(int i, real_t &r)
{
const real_t diff = fabs(exact_d[i] - gridf_d[i]);
r += w_d[i%NQ] * detJ_d[i] * pow(diff, p);
}, SumReducer<real_t> {}, true, workspace);
error = pow(error, 1./p);
}
else
{
mfem::reduce(q_vals.Size(), error,
[=] MFEM_HOST_DEVICE(int i, real_t &r)
{
const real_t diff = fabs(exact_d[i] - gridf_d[i]);
r = fmax(r, diff);
}, MaxReducer<real_t> {}, true, workspace);
}
return error;
}
const FiniteElement *fe;
ElementTransformation *T;
Vector vals;
@@ -4453,7 +4334,7 @@ real_t LSZZErrorEstimator(BilinearFormIntegrator &blfi, // input
u.GetSubVector(udofs, ul);
utrans.InvTransformPrimal(ul);
Transf = ufes->GetElementTransformation(ielem);
const auto *dummy = ufes->GetFE(ielem);
FiniteElement *dummy = nullptr;
blfi.ComputeElementFlux(*ufes->GetFE(ielem), *Transf, ul,
*dummy, fl, with_coeff, ir);
+31 -34
View File
@@ -161,8 +161,7 @@ static void EADiffusionAssemble3D(const int NE,
auto B = Reshape(b.Read(), Q1D, D1D);
auto G = Reshape(g.Read(), Q1D, D1D);
auto D = Reshape(padata.Read(), Q1D, Q1D, Q1D, 6, NE);
auto A = Reshape(add ? eadata.ReadWrite() : eadata.Write(),
D1D, D1D, D1D, D1D, D1D, D1D, NE);
auto A = Reshape(eadata.ReadWrite(), D1D, D1D, D1D, D1D, D1D, D1D, NE);
mfem::forall_3D(NE, D1D, D1D, D1D, [=] MFEM_HOST_DEVICE (int e)
{
const int D1D = T_D1D ? T_D1D : d1d;
@@ -247,60 +246,58 @@ void DiffusionIntegrator::AssembleEA(const FiniteElementSpace &fes,
Vector &ea_data,
const bool add)
{
MFEM_PERF_FUNCTION;
AssemblePA(fes);
ne = fes.GetMesh()->GetNE();
const Array<real_t> &B = maps->B;
const Array<real_t> &G = maps->G;
decltype(&EADiffusionAssemble1D<>) kernel = nullptr;
if (dim == 1)
{
switch ((dofs1D << 4 ) | quad1D)
{
case 0x22: kernel = EADiffusionAssemble1D<2,2>;
case 0x33: kernel = EADiffusionAssemble1D<3,3>;
case 0x44: kernel = EADiffusionAssemble1D<4,4>;
case 0x55: kernel = EADiffusionAssemble1D<5,5>;
case 0x66: kernel = EADiffusionAssemble1D<6,6>;
case 0x77: kernel = EADiffusionAssemble1D<7,7>;
case 0x88: kernel = EADiffusionAssemble1D<8,8>;
case 0x99: kernel = EADiffusionAssemble1D<9,9>;
default: kernel = EADiffusionAssemble1D<>;
case 0x22: return EADiffusionAssemble1D<2,2>(ne,B,G,pa_data,ea_data,add);
case 0x33: return EADiffusionAssemble1D<3,3>(ne,B,G,pa_data,ea_data,add);
case 0x44: return EADiffusionAssemble1D<4,4>(ne,B,G,pa_data,ea_data,add);
case 0x55: return EADiffusionAssemble1D<5,5>(ne,B,G,pa_data,ea_data,add);
case 0x66: return EADiffusionAssemble1D<6,6>(ne,B,G,pa_data,ea_data,add);
case 0x77: return EADiffusionAssemble1D<7,7>(ne,B,G,pa_data,ea_data,add);
case 0x88: return EADiffusionAssemble1D<8,8>(ne,B,G,pa_data,ea_data,add);
case 0x99: return EADiffusionAssemble1D<9,9>(ne,B,G,pa_data,ea_data,add);
default: return EADiffusionAssemble1D(ne,B,G,pa_data,ea_data,add,
dofs1D,quad1D);
}
}
else if (dim == 2)
{
switch ((dofs1D << 4 ) | quad1D)
{
case 0x22: kernel = EADiffusionAssemble2D<2,2>;
case 0x33: kernel = EADiffusionAssemble2D<3,3>;
case 0x44: kernel = EADiffusionAssemble2D<4,4>;
case 0x55: kernel = EADiffusionAssemble2D<5,5>;
case 0x66: kernel = EADiffusionAssemble2D<6,6>;
case 0x77: kernel = EADiffusionAssemble2D<7,7>;
case 0x88: kernel = EADiffusionAssemble2D<8,8>;
case 0x99: kernel = EADiffusionAssemble2D<9,9>;
default: kernel = EADiffusionAssemble2D<>;
case 0x22: return EADiffusionAssemble2D<2,2>(ne,B,G,pa_data,ea_data,add);
case 0x33: return EADiffusionAssemble2D<3,3>(ne,B,G,pa_data,ea_data,add);
case 0x44: return EADiffusionAssemble2D<4,4>(ne,B,G,pa_data,ea_data,add);
case 0x55: return EADiffusionAssemble2D<5,5>(ne,B,G,pa_data,ea_data,add);
case 0x66: return EADiffusionAssemble2D<6,6>(ne,B,G,pa_data,ea_data,add);
case 0x77: return EADiffusionAssemble2D<7,7>(ne,B,G,pa_data,ea_data,add);
case 0x88: return EADiffusionAssemble2D<8,8>(ne,B,G,pa_data,ea_data,add);
case 0x99: return EADiffusionAssemble2D<9,9>(ne,B,G,pa_data,ea_data,add);
default: return EADiffusionAssemble2D(ne,B,G,pa_data,ea_data,add,
dofs1D,quad1D);
}
}
else if (dim == 3)
{
switch ((dofs1D << 4 ) | quad1D)
{
case 0x23: kernel = EADiffusionAssemble3D<2,3>;
case 0x34: kernel = EADiffusionAssemble3D<3,4>;
case 0x45: kernel = EADiffusionAssemble3D<4,5>;
case 0x56: kernel = EADiffusionAssemble3D<5,6>;
case 0x67: kernel = EADiffusionAssemble3D<6,7>;
case 0x78: kernel = EADiffusionAssemble3D<7,8>;
case 0x89: kernel = EADiffusionAssemble3D<8,9>;
default: kernel = EADiffusionAssemble3D<>;
case 0x23: return EADiffusionAssemble3D<2,3>(ne,B,G,pa_data,ea_data,add);
case 0x34: return EADiffusionAssemble3D<3,4>(ne,B,G,pa_data,ea_data,add);
case 0x45: return EADiffusionAssemble3D<4,5>(ne,B,G,pa_data,ea_data,add);
case 0x56: return EADiffusionAssemble3D<5,6>(ne,B,G,pa_data,ea_data,add);
case 0x67: return EADiffusionAssemble3D<6,7>(ne,B,G,pa_data,ea_data,add);
case 0x78: return EADiffusionAssemble3D<7,8>(ne,B,G,pa_data,ea_data,add);
case 0x89: return EADiffusionAssemble3D<8,9>(ne,B,G,pa_data,ea_data,add);
default: return EADiffusionAssemble3D(ne,B,G,pa_data,ea_data,add,
dofs1D,quad1D);
}
}
MFEM_VERIFY(kernel != nullptr, "Unknown kernel.");
kernel(ne,B,G,pa_data,ea_data,add,dofs1D,quad1D);
// Free the PA data:
pa_data.Destroy();
MFEM_ABORT("Unknown kernel.");
}
}
+25 -26
View File
@@ -1009,7 +1009,6 @@ inline void SmemPADiffusionApply3D(const int NE,
auto d = Reshape(d_.Read(), Q1D, Q1D, Q1D, symmetric ? 6 : 9, NE);
auto x = Reshape(x_.Read(), D1D, D1D, D1D, NE);
auto y = Reshape(y_.ReadWrite(), D1D, D1D, D1D, NE);
MFEM_VERIFY(D1D <= Q1D, "THREAD_DIRECT requires D1D <= Q1D");
mfem::forall_3D(NE, Q1D, Q1D, Q1D, [=] MFEM_HOST_DEVICE (int e)
{
const int D1D = T_D1D ? T_D1D : d1d;
@@ -1039,11 +1038,11 @@ inline void SmemPADiffusionApply3D(const int NE,
real_t (*QDD0)[MD1][MD1] = (real_t (*)[MD1][MD1]) (sm0+0);
real_t (*QDD1)[MD1][MD1] = (real_t (*)[MD1][MD1]) (sm0+1);
real_t (*QDD2)[MD1][MD1] = (real_t (*)[MD1][MD1]) (sm0+2);
MFEM_FOREACH_THREAD_DIRECT(dz,z,D1D)
MFEM_FOREACH_THREAD(dz,z,D1D)
{
MFEM_FOREACH_THREAD_DIRECT(dy,y,D1D)
MFEM_FOREACH_THREAD(dy,y,D1D)
{
MFEM_FOREACH_THREAD_DIRECT(dx,x,D1D)
MFEM_FOREACH_THREAD(dx,x,D1D)
{
X[dz][dy][dx] = x(dx,dy,dz,e);
}
@@ -1051,9 +1050,9 @@ inline void SmemPADiffusionApply3D(const int NE,
}
if (MFEM_THREAD_ID(z) == 0)
{
MFEM_FOREACH_THREAD_DIRECT(dy,y,D1D)
MFEM_FOREACH_THREAD(dy,y,D1D)
{
MFEM_FOREACH_THREAD_DIRECT(qx,x,Q1D)
MFEM_FOREACH_THREAD(qx,x,Q1D)
{
B[qx][dy] = b(qx,dy);
G[qx][dy] = g(qx,dy);
@@ -1061,11 +1060,11 @@ inline void SmemPADiffusionApply3D(const int NE,
}
}
MFEM_SYNC_THREAD;
MFEM_FOREACH_THREAD_DIRECT(dz,z,D1D)
MFEM_FOREACH_THREAD(dz,z,D1D)
{
MFEM_FOREACH_THREAD_DIRECT(dy,y,D1D)
MFEM_FOREACH_THREAD(dy,y,D1D)
{
MFEM_FOREACH_THREAD_DIRECT(qx,x,Q1D)
MFEM_FOREACH_THREAD(qx,x,Q1D)
{
real_t u = 0.0, v = 0.0;
MFEM_UNROLL(MD1)
@@ -1081,11 +1080,11 @@ inline void SmemPADiffusionApply3D(const int NE,
}
}
MFEM_SYNC_THREAD;
MFEM_FOREACH_THREAD_DIRECT(dz,z,D1D)
MFEM_FOREACH_THREAD(dz,z,D1D)
{
MFEM_FOREACH_THREAD_DIRECT(qy,y,Q1D)
MFEM_FOREACH_THREAD(qy,y,Q1D)
{
MFEM_FOREACH_THREAD_DIRECT(qx,x,Q1D)
MFEM_FOREACH_THREAD(qx,x,Q1D)
{
real_t u = 0.0, v = 0.0, w = 0.0;
MFEM_UNROLL(MD1)
@@ -1102,11 +1101,11 @@ inline void SmemPADiffusionApply3D(const int NE,
}
}
MFEM_SYNC_THREAD;
MFEM_FOREACH_THREAD_DIRECT(qz,z,Q1D)
MFEM_FOREACH_THREAD(qz,z,Q1D)
{
MFEM_FOREACH_THREAD_DIRECT(qy,y,Q1D)
MFEM_FOREACH_THREAD(qy,y,Q1D)
{
MFEM_FOREACH_THREAD_DIRECT(qx,x,Q1D)
MFEM_FOREACH_THREAD(qx,x,Q1D)
{
real_t u = 0.0, v = 0.0, w = 0.0;
MFEM_UNROLL(MD1)
@@ -1137,9 +1136,9 @@ inline void SmemPADiffusionApply3D(const int NE,
MFEM_SYNC_THREAD;
if (MFEM_THREAD_ID(z) == 0)
{
MFEM_FOREACH_THREAD_DIRECT(dy,y,D1D)
MFEM_FOREACH_THREAD(dy,y,D1D)
{
MFEM_FOREACH_THREAD_DIRECT(qx,x,Q1D)
MFEM_FOREACH_THREAD(qx,x,Q1D)
{
Bt[dy][qx] = b(qx,dy);
Gt[dy][qx] = g(qx,dy);
@@ -1147,11 +1146,11 @@ inline void SmemPADiffusionApply3D(const int NE,
}
}
MFEM_SYNC_THREAD;
MFEM_FOREACH_THREAD_DIRECT(qz,z,Q1D)
MFEM_FOREACH_THREAD(qz,z,Q1D)
{
MFEM_FOREACH_THREAD_DIRECT(qy,y,Q1D)
MFEM_FOREACH_THREAD(qy,y,Q1D)
{
MFEM_FOREACH_THREAD_DIRECT(dx,x,D1D)
MFEM_FOREACH_THREAD(dx,x,D1D)
{
real_t u = 0.0, v = 0.0, w = 0.0;
MFEM_UNROLL(MQ1)
@@ -1168,11 +1167,11 @@ inline void SmemPADiffusionApply3D(const int NE,
}
}
MFEM_SYNC_THREAD;
MFEM_FOREACH_THREAD_DIRECT(qz,z,Q1D)
MFEM_FOREACH_THREAD(qz,z,Q1D)
{
MFEM_FOREACH_THREAD_DIRECT(dy,y,D1D)
MFEM_FOREACH_THREAD(dy,y,D1D)
{
MFEM_FOREACH_THREAD_DIRECT(dx,x,D1D)
MFEM_FOREACH_THREAD(dx,x,D1D)
{
real_t u = 0.0, v = 0.0, w = 0.0;
MFEM_UNROLL(Q1D)
@@ -1189,11 +1188,11 @@ inline void SmemPADiffusionApply3D(const int NE,
}
}
MFEM_SYNC_THREAD;
MFEM_FOREACH_THREAD_DIRECT(dz,z,D1D)
MFEM_FOREACH_THREAD(dz,z,D1D)
{
MFEM_FOREACH_THREAD_DIRECT(dy,y,D1D)
MFEM_FOREACH_THREAD(dy,y,D1D)
{
MFEM_FOREACH_THREAD_DIRECT(dx,x,D1D)
MFEM_FOREACH_THREAD(dx,x,D1D)
{
real_t u = 0.0, v = 0.0, w = 0.0;
MFEM_UNROLL(MQ1)
-4
View File
@@ -39,8 +39,6 @@ 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);
@@ -90,8 +88,6 @@ 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
+56 -29
View File
@@ -23,8 +23,6 @@ namespace mfem
void MassIntegrator::AssemblePA(const FiniteElementSpace &fes)
{
MFEM_PERF_FUNCTION;
const MemoryType mt = (pa_mt == MemoryType::DEFAULT) ?
Device::GetDeviceMemoryType() : pa_mt;
@@ -61,23 +59,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 = 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)
@@ -108,22 +109,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)
@@ -141,8 +170,6 @@ void MassIntegrator::AssembleDiagonalPA(Vector &diag)
void MassIntegrator::AddMultPA(const Vector &x, Vector &y) const
{
MFEM_PERF_FUNCTION;
if (DeviceCanUseCeed())
{
ceedOp->AddMult(x, y);
-2
View File
@@ -242,8 +242,6 @@ void DomainLFIntegrator::AssembleDevice(const FiniteElementSpace &fes,
const Array<int> &markers,
Vector &b)
{
MFEM_PERF_FUNCTION;
const FiniteElement &fe = *fes.GetTypicalFE();
const int qorder = oa * fe.GetOrder() + ob;
const Geometry::Type gtype = fe.GetGeomType();
+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));
-3
View File
@@ -161,7 +161,6 @@ bool LinearForm::SupportsDevice() const
void LinearForm::UseFastAssembly(bool use_fa)
{
MFEM_PERF_FUNCTION;
fast_assembly = use_fa;
if (fast_assembly && SupportsDevice() && !ext)
@@ -172,8 +171,6 @@ void LinearForm::UseFastAssembly(bool use_fa)
void LinearForm::Assemble()
{
MFEM_PERF_FUNCTION;
Array<int> vdofs;
ElementTransformation *eltrans;
Vector elemvect;
+33 -19
View File
@@ -15,16 +15,10 @@
namespace mfem
{
LinearFormExtension::LinearFormExtension(LinearForm *lf): lf(lf)
{
MFEM_PERF_FUNCTION;
Update();
}
LinearFormExtension::LinearFormExtension(LinearForm *lf): lf(lf) { Update(); }
void LinearFormExtension::Assemble()
{
MFEM_PERF_FUNCTION;
const FiniteElementSpace &fes = *lf->FESpace();
MFEM_VERIFY(lf->SupportsDevice(), "Not supported.");
MFEM_VERIFY(lf->Size() == fes.GetVSize(), "LinearForm size does not "
@@ -57,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)
@@ -99,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;
});
}
@@ -119,8 +112,6 @@ void LinearFormExtension::Assemble()
void LinearFormExtension::Update()
{
MFEM_PERF_FUNCTION;
const FiniteElementSpace &fes = *lf->FESpace();
const Mesh &mesh = *fes.GetMesh();
constexpr ElementDofOrdering ordering = ElementDofOrdering::LEXICOGRAPHIC;
@@ -134,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");
@@ -144,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,
+1 -3
View File
@@ -14,7 +14,6 @@
#include "../general/array.hpp"
#include "../linalg/vector.hpp"
#include "fespace.hpp"
namespace mfem
{
@@ -26,8 +25,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;
+1 -1
View File
@@ -673,7 +673,7 @@ public:
int myid;
MPI_Comm_rank(comm, &myid);
int seed = (seed_ > 0) ? seed_ + myid : time(nullptr) + myid;
int seed = (seed_ > 0) ? seed_ + myid : (int)time(0) + myid;
SetSeed(seed);
}
#else
-2
View File
@@ -365,8 +365,6 @@ 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,8 +360,6 @@ 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
@@ -419,8 +417,6 @@ 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,8 +184,6 @@ 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;
+54 -139
View File
@@ -10,7 +10,6 @@
// CONTRIBUTING.md for details.
#include "multigrid.hpp"
#include "../general/annotation.hpp"
namespace mfem
{
@@ -18,10 +17,7 @@ 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_,
@@ -33,18 +29,12 @@ 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])
@@ -66,17 +56,16 @@ void MultigridBase::InitVectors() const
X.SetSize(M, nrhs);
Y.SetSize(M, nrhs);
R.SetSize(M, nrhs);
for (int i = 0; i < M; ++i)
Z.SetSize(M, nrhs);
for (int i = 0; i < X.NumRows(); ++i)
{
const int n = operators[i]->Height();
for (int j = 0; j < nrhs; ++j)
for (int j = 0; j < X.NumCols(); ++j)
{
if (i < M - 1)
{
X(i, j) = new Vector(n);
Y(i, j) = new Vector(n);
}
X(i, j) = new Vector(n);
Y(i, j) = new Vector(n);
R(i, j) = new Vector(n);
Z(i, j) = new Vector(n);
}
}
}
@@ -87,12 +76,10 @@ void MultigridBase::EraseVectors() const
{
for (int j = 0; j < X.NumCols(); ++j)
{
if (i < X.NumRows() - 1)
{
delete X(i, j);
delete Y(i, j);
}
delete X(i, j);
delete Y(i, j);
delete R(i, j);
delete Z(i, j);
}
}
}
@@ -108,12 +95,6 @@ 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_)
{
@@ -124,24 +105,25 @@ void MultigridBase::SetCycleType(CycleType cycleType_, int preSmoothingSteps_,
void MultigridBase::Mult(const Vector& x, Vector& y) const
{
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
Array<const Vector*> X_(1);
Array<Vector*> Y_(1);
X_[0] = &x;
Y_[0] = &y;
ArrayMult(X_, Y_);
}
void MultigridBase::ArrayMult(const Array<const Vector*>& X_,
Array<Vector*>& Y_) const
{
MFEM_PERF_FUNCTION;
MFEM_ASSERT(operators.Size() > 0,
"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();
@@ -152,163 +134,96 @@ 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) = const_cast<Vector*>(X_[j]);
Y(M - 1, j) = Y_[j];
*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);
}
const bool zero = !iterative_mode;
Cycle(M - 1, zero);
}
void MultigridBase::SmoothingStep(int level, bool zero, bool transpose) const
{
MFEM_PERF_FUNCTION;
// 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)
{
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_);
Array<Vector *> X_(X[level], nrhs), Y_(Y[level], nrhs);
GetSmootherAtLevel(level)->ArrayMult(X_, Y_);
}
else
{
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_);
Array<Vector *> Y_(Y[level], nrhs), R_(R[level], nrhs),
Z_(Z[level], nrhs);
for (int j = 0; j < nrhs; ++j)
{
// *R_[j] = *X(level, j) - *R_[j]
subtract(*X(level, j), *R_[j], *R_[j]);
*R_[j] = *X(level, j);
}
GetOperatorAtLevel(level)->ArrayAddMult(Y_, R_, -1.0);
if (transpose)
{
GetSmootherAtLevel(level)->ArrayAddMultTranspose(cR_, Y_);
GetSmootherAtLevel(level)->ArrayMultTranspose(R_, Z_);
}
else
{
GetSmootherAtLevel(level)->ArrayAddMult(cR_, Y_);
GetSmootherAtLevel(level)->ArrayMult(R_, Z_);
}
}
}
void MultigridBase::CoarseSolve(bool zero) const
{
MFEM_PERF_FUNCTION;
// See the comment about iterative mode in SmoothingStep()
coarse_solver->iterative_mode = false;
if (zero)
{
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]);
*Y_[j] += *Z_[j];
}
coarse_solver->ArrayAddMult(cR_, Y_);
}
}
void MultigridBase::Cycle(int level, bool zero) const
void MultigridBase::Cycle(int level) 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)
if (level == 0)
{
SmoothingStep(0, zero, false);
SmoothingStep(0, true, false);
return;
}
// Pre-smooth
for (int i = 0; i < preSmoothingSteps; ++i)
{
SmoothingStep(level, zero && (i == 0), false);
}
// Coarse solve with 'coarse_solver'
if (level == 0)
{
CoarseSolve(preSmoothingSteps == 0 && zero);
goto mg_post_smooth;
SmoothingStep(level, (cycleType == CycleType::VCYCLE && i == 0), false);
}
// Compute residual and restrict
if (preSmoothingSteps == 0 && zero)
{
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_);
Array<Vector *> Y_(Y[level], nrhs), R_(R[level], nrhs),
X_(X[level - 1], nrhs);
for (int j = 0; j < nrhs; ++j)
{
// *R_[j] = *X(level, j) - *R_[j]
subtract(*X(level, j), *R_[j], *R_[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;
}
GetProlongationAtLevel(level - 1)->ArrayMultTranspose(cR_, X_);
}
// Corrections
Cycle(level - 1, true);
Cycle(level - 1);
if (cycleType == CycleType::WCYCLE)
{
// 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); }
Cycle(level - 1);
}
// Prolongate and add
{
const Array<const Vector *> cY_lm1((const Vector **)(Y[level - 1]), nrhs);
Array<Vector *> Y_l(Y[level], nrhs);
if (preSmoothingSteps == 0 && zero)
Array<Vector *> Y_(Y[level - 1], nrhs), Z_(Z[level], nrhs);
GetProlongationAtLevel(level - 1)->ArrayMult(Y_, Z_);
for (int j = 0; j < nrhs; ++j)
{
GetProlongationAtLevel(level - 1)->ArrayMult(cY_lm1, Y_l);
}
else
{
GetProlongationAtLevel(level - 1)->ArrayAddMult(cY_lm1, Y_l);
*Y(level, j) += *Z_[j];
}
}
mg_post_smooth:
// Post-smooth
for (int i = 0; i < postSmoothingSteps; ++i)
{
+2 -20
View File
@@ -36,14 +36,12 @@ 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;
mutable Array2D<Vector*> X, Y, R, Z;
mutable int nrhs;
public:
@@ -67,16 +65,6 @@ 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(); }
@@ -130,14 +118,11 @@ public:
private:
/// Application of a multigrid cycle at particular level
void Cycle(int level, bool zero) const;
void Cycle(int level) 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;
@@ -217,9 +202,6 @@ 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,8 +124,6 @@ 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
+35 -60
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)
@@ -63,11 +62,9 @@ ParFiniteElementSpace::ParFiniteElementSpace(
ParFiniteElementSpace::ParFiniteElementSpace(
ParMesh *pm, const FiniteElementCollection *f, int dim, int ordering)
: FiniteElementSpace((MFEM_PERF_BEGIN(_MFEM_FUNC_NAME), pm),
f, dim, ordering)
: FiniteElementSpace(pm, f, dim, ordering)
{
ParInit(pm);
MFEM_PERF_END(_MFEM_FUNC_NAME);
}
ParFiniteElementSpace::ParFiniteElementSpace(
@@ -94,7 +91,6 @@ ParNURBSExtension *ParFiniteElementSpace::MakeLocalNURBSext(
void ParFiniteElementSpace::ParInit(ParMesh *pm)
{
MFEM_PERF_FUNCTION;
pmesh = pm;
pncmesh = nullptr;
@@ -184,7 +180,6 @@ void ParFiniteElementSpace::CommunicateGhostOrder()
void ParFiniteElementSpace::Construct()
{
MFEM_PERF_FUNCTION;
if (NURBSext)
{
ConstructTrueNURBSDofs();
@@ -843,8 +838,6 @@ void ParFiniteElementSpace::Build_Dof_TrueDof_Matrix() const // matrix P
if (P) { return; }
MFEM_PERF_FUNCTION;
if (!nd_strias)
{
// Safe to assume 1-1 correspondence between shared dofs
@@ -1430,7 +1423,6 @@ const Operator *ParFiniteElementSpace::GetRestrictionOperator() const
if (NRanks == 1)
{
R_transpose.reset(new IdentityOperator(GetTrueVSize()));
Rconf = new IdentityOperator(GetTrueVSize());
}
else
{
@@ -1443,8 +1435,8 @@ const Operator *ParFiniteElementSpace::GetRestrictionOperator() const
R_transpose.reset(
new DeviceConformingProlongationOperator(*this, true));
}
Rconf = new TransposeOperator(*R_transpose);
}
Rconf = new TransposeOperator(*R_transpose);
return Rconf;
}
else
@@ -4495,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,
@@ -4537,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];
@@ -4563,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];
@@ -4660,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);
@@ -4681,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++)
{
@@ -4715,6 +4694,7 @@ ParFiniteElementSpace::ParallelDerefinementMatrix(int old_ndofs,
}
}
messages.clear();
offd->Finalize(0);
offd->SetWidth(static_cast<int>(col_map.size()));
@@ -4966,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);
@@ -5284,7 +5259,7 @@ DeviceConformingProlongationOperator::DeviceConformingProlongationOperator(
gc.GetNeighborLTDofTable(nbr_ltdof);
const int nb_connections = nbr_ltdof.Size_of_connections();
shr_ltdof.SetSize(nb_connections);
if (nb_connections > 0) { shr_ltdof.CopyFrom(nbr_ltdof.GetJ()); }
shr_ltdof.CopyFrom(nbr_ltdof.GetJ());
shr_buf.SetSize(nb_connections);
shr_buf.UseDevice(true);
shr_buf_offsets = nbr_ltdof.GetIMemory();
@@ -5313,7 +5288,7 @@ DeviceConformingProlongationOperator::DeviceConformingProlongationOperator(
gc.GetNeighborLDofTable(nbr_ldof);
const int nb_connections = nbr_ldof.Size_of_connections();
ext_ldof.SetSize(nb_connections);
if (nb_connections > 0) { ext_ldof.CopyFrom(nbr_ldof.GetJ()); }
ext_ldof.CopyFrom(nbr_ldof.GetJ());
ext_ldof.GetMemory().UseDevice(true);
ext_buf.SetSize(nb_connections);
ext_buf.UseDevice(true);
-3
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;
-2
View File
@@ -45,8 +45,6 @@ void ParLinearForm::MakeRef(ParFiniteElementSpace *pf, Vector &v, int v_offset)
void ParLinearForm::Assemble()
{
MFEM_PERF_FUNCTION;
LinearForm::Assemble();
if (interior_face_integs.Size())
+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;
+280 -3
View File
@@ -9,16 +9,278 @@
// terms of the BSD-3 license. We welcome feedback and contributions, see file
// CONTRIBUTING.md for details.
#include "det.hpp"
#include "../quadinterpolator.hpp"
#include "../../general/forall.hpp"
#include "../../linalg/dtensor.hpp"
#include "../../fem/kernels.hpp"
#include "../../linalg/kernels.hpp"
using namespace mfem;
namespace mfem
{
namespace internal
{
namespace quadrature_interpolator
{
static void Det1D(const int NE,
const real_t *b,
const real_t *g,
const real_t *x,
real_t *y,
const int d1d,
const int q1d,
Vector *d_buff = nullptr)
{
MFEM_CONTRACT_VAR(b);
MFEM_CONTRACT_VAR(d_buff);
const auto G = Reshape(g, q1d, d1d);
const auto X = Reshape(x, d1d, NE);
auto Y = Reshape(y, q1d, NE);
mfem::forall(NE, [=] MFEM_HOST_DEVICE (int e)
{
for (int q = 0; q < q1d; q++)
{
real_t u = 0.0;
for (int d = 0; d < d1d; d++)
{
u += G(q, d) * X(d, e);
}
Y(q, e) = u;
}
});
}
template<int T_D1D = 0, int T_Q1D = 0>
static void Det2D(const int NE,
const real_t *b,
const real_t *g,
const real_t *x,
real_t *y,
const int d1d = 0,
const int q1d = 0,
Vector *d_buff = nullptr)
{
MFEM_CONTRACT_VAR(d_buff);
static constexpr int SDIM = 2;
static constexpr int NBZ = 1;
const int D1D = T_D1D ? T_D1D : d1d;
const int Q1D = T_Q1D ? T_Q1D : q1d;
const auto B = Reshape(b, Q1D, D1D);
const auto G = Reshape(g, Q1D, D1D);
const auto X = Reshape(x, D1D, D1D, SDIM, NE);
auto Y = Reshape(y, Q1D, Q1D, NE);
mfem::forall_2D_batch(NE, Q1D, Q1D, NBZ, [=] MFEM_HOST_DEVICE (int e)
{
constexpr int MQ1 = T_Q1D ? T_Q1D : DofQuadLimits::MAX_Q1D;
constexpr int MD1 = T_D1D ? T_D1D : DofQuadLimits::MAX_D1D;
const int D1D = T_D1D ? T_D1D : d1d;
const int Q1D = T_Q1D ? T_Q1D : q1d;
MFEM_SHARED real_t BG[2][MQ1*MD1];
MFEM_SHARED real_t XY[SDIM][NBZ][MD1*MD1];
MFEM_SHARED real_t DQ[2*SDIM][NBZ][MD1*MQ1];
MFEM_SHARED real_t QQ[2*SDIM][NBZ][MQ1*MQ1];
kernels::internal::LoadX<MD1,NBZ>(e,D1D,X,XY);
kernels::internal::LoadBG<MD1,MQ1>(D1D,Q1D,B,G,BG);
kernels::internal::GradX<MD1,MQ1,NBZ>(D1D,Q1D,BG,XY,DQ);
kernels::internal::GradY<MD1,MQ1,NBZ>(D1D,Q1D,BG,DQ,QQ);
MFEM_FOREACH_THREAD(qy,y,Q1D)
{
MFEM_FOREACH_THREAD(qx,x,Q1D)
{
real_t J[4];
kernels::internal::PullGrad<MQ1,NBZ>(Q1D,qx,qy,QQ,J);
Y(qx,qy,e) = kernels::Det<2>(J);
}
}
});
}
template<int T_D1D = 0, int T_Q1D = 0>
static void Det2DSurface(const int NE,
const real_t *b,
const real_t *g,
const real_t *x,
real_t *y,
const int d1d = 0,
const int q1d = 0,
Vector *d_buff = nullptr)
{
MFEM_CONTRACT_VAR(d_buff);
static constexpr int SDIM = 3;
static constexpr int NBZ = 1;
const int D1D = T_D1D ? T_D1D : d1d;
const int Q1D = T_Q1D ? T_Q1D : q1d;
const auto B = Reshape(b, Q1D, D1D);
const auto G = Reshape(g, Q1D, D1D);
const auto X = Reshape(x, D1D, D1D, SDIM, NE);
auto Y = Reshape(y, Q1D, Q1D, NE);
mfem::forall_2D_batch(NE, Q1D, Q1D, NBZ, [=] MFEM_HOST_DEVICE (int e)
{
constexpr int MQ1 = T_Q1D ? T_Q1D : DofQuadLimits::MAX_Q1D;
constexpr int MD1 = T_D1D ? T_D1D : DofQuadLimits::MAX_D1D;
const int D1D = T_D1D ? T_D1D : d1d;
const int Q1D = T_Q1D ? T_Q1D : q1d;
const int tidz = MFEM_THREAD_ID(z);
MFEM_SHARED real_t BG[2][MQ1*MD1];
MFEM_SHARED real_t XYZ[SDIM][NBZ][MD1*MD1];
MFEM_SHARED real_t DQ[2*SDIM][NBZ][MD1*MQ1];
kernels::internal::LoadBG<MD1,MQ1>(D1D,Q1D,B,G,BG);
// Load XYZ components
MFEM_FOREACH_THREAD(dy,y,D1D)
{
MFEM_FOREACH_THREAD(dx,x,D1D)
{
for (int d = 0; d < SDIM; ++d)
{
XYZ[d][tidz][dx + dy*D1D] = X(dx,dy,d,e);
}
}
}
MFEM_SYNC_THREAD;
ConstDeviceMatrix B_mat(BG[0], D1D, Q1D);
ConstDeviceMatrix G_mat(BG[1], D1D, Q1D);
// x contraction
MFEM_FOREACH_THREAD(dy,y,D1D)
{
MFEM_FOREACH_THREAD(qx,x,Q1D)
{
for (int d = 0; d < SDIM; ++d)
{
real_t u = 0.0;
real_t v = 0.0;
for (int dx = 0; dx < D1D; ++dx)
{
const real_t xval = XYZ[d][tidz][dx + dy*D1D];
u += xval * G_mat(dx,qx);
v += xval * B_mat(dx,qx);
}
DQ[d][tidz][dy + qx*D1D] = u;
DQ[3 + d][tidz][dy + qx*D1D] = v;
}
}
}
MFEM_SYNC_THREAD;
// y contraction and determinant computation
MFEM_FOREACH_THREAD(qy,y,Q1D)
{
MFEM_FOREACH_THREAD(qx,x,Q1D)
{
real_t J_[6] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0};
for (int d = 0; d < SDIM; ++d)
{
for (int dy = 0; dy < D1D; ++dy)
{
J_[d] += DQ[d][tidz][dy + qx*D1D] * B_mat(dy,qy);
J_[3 + d] += DQ[3 + d][tidz][dy + qx*D1D] * G_mat(dy,qy);
}
}
DeviceTensor<2> J(J_, 3, 2);
const real_t E = J(0,0)*J(0,0) + J(1,0)*J(1,0) + J(2,0)*J(2,0);
const real_t F = J(0,0)*J(0,1) + J(1,0)*J(1,1) + J(2,0)*J(2,1);
const real_t G = J(0,1)*J(0,1) + J(1,1)*J(1,1) + J(2,1)*J(2,1);
Y(qx,qy,e) = std::sqrt(E*G - F*F);
}
}
});
}
template<int T_D1D = 0, int T_Q1D = 0, bool SMEM = true>
static void Det3D(const int NE,
const real_t *b,
const real_t *g,
const real_t *x,
real_t *y,
const int d1d = 0,
const int q1d = 0,
Vector *d_buff = nullptr) // used only with SMEM = false
{
constexpr int DIM = 3;
static constexpr int GRID = SMEM ? 0 : 128;
const int D1D = T_D1D ? T_D1D : d1d;
const int Q1D = T_Q1D ? T_Q1D : q1d;
const auto B = Reshape(b, Q1D, D1D);
const auto G = Reshape(g, Q1D, D1D);
const auto X = Reshape(x, D1D, D1D, D1D, DIM, NE);
auto Y = Reshape(y, Q1D, Q1D, Q1D, NE);
real_t *GM = nullptr;
if (!SMEM)
{
const DeviceDofQuadLimits &limits = DeviceDofQuadLimits::Get();
const int max_q1d = T_Q1D ? T_Q1D : limits.MAX_Q1D;
const int max_d1d = T_D1D ? T_D1D : limits.MAX_D1D;
const int max_qd = std::max(max_q1d, max_d1d);
const int mem_size = max_qd * max_qd * max_qd * 9;
d_buff->SetSize(2*mem_size*GRID);
GM = d_buff->Write();
}
mfem::forall_3D_grid(NE, Q1D, Q1D, Q1D, GRID, [=] MFEM_HOST_DEVICE (int e)
{
static constexpr int MQ1 = T_Q1D ? T_Q1D :
(SMEM ? DofQuadLimits::MAX_DET_1D : DofQuadLimits::MAX_Q1D);
static constexpr int MD1 = T_D1D ? T_D1D :
(SMEM ? DofQuadLimits::MAX_DET_1D : DofQuadLimits::MAX_D1D);
static constexpr int MDQ = MQ1 > MD1 ? MQ1 : MD1;
static constexpr int MSZ = MDQ * MDQ * MDQ * 9;
const int bid = MFEM_BLOCK_ID(x);
MFEM_SHARED real_t BG[2][MQ1*MD1];
MFEM_SHARED real_t SM0[SMEM?MSZ:1];
MFEM_SHARED real_t SM1[SMEM?MSZ:1];
real_t *lm0 = SMEM ? SM0 : GM + MSZ*bid;
real_t *lm1 = SMEM ? SM1 : GM + MSZ*(GRID+bid);
real_t (*DDD)[MD1*MD1*MD1] = (real_t (*)[MD1*MD1*MD1]) (lm0);
real_t (*DDQ)[MD1*MD1*MQ1] = (real_t (*)[MD1*MD1*MQ1]) (lm1);
real_t (*DQQ)[MD1*MQ1*MQ1] = (real_t (*)[MD1*MQ1*MQ1]) (lm0);
real_t (*QQQ)[MQ1*MQ1*MQ1] = (real_t (*)[MQ1*MQ1*MQ1]) (lm1);
kernels::internal::LoadX<MD1>(e,D1D,X,DDD);
kernels::internal::LoadBG<MD1,MQ1>(D1D,Q1D,B,G,BG);
kernels::internal::GradX<MD1,MQ1>(D1D,Q1D,BG,DDD,DDQ);
kernels::internal::GradY<MD1,MQ1>(D1D,Q1D,BG,DDQ,DQQ);
kernels::internal::GradZ<MD1,MQ1>(D1D,Q1D,BG,DQQ,QQQ);
MFEM_FOREACH_THREAD(qz,z,Q1D)
{
MFEM_FOREACH_THREAD(qy,y,Q1D)
{
MFEM_FOREACH_THREAD(qx,x,Q1D)
{
real_t J[9];
kernels::internal::PullGrad<MQ1>(Q1D, qx,qy,qz, QQQ, J);
Y(qx,qy,qz,e) = kernels::Det<3>(J);
}
}
}
});
}
void InitDetKernels()
{
using k = QuadratureInterpolator::DetKernels;
@@ -40,12 +302,27 @@ void InitDetKernels()
}
} // namespace quadrature_interpolator
} // namespace internal
/// @cond Suppress_Doxygen_warnings
QuadratureInterpolator::DetKernelType
QuadratureInterpolator::DetKernels::Fallback(
namespace
{
using DetKernel = QuadratureInterpolator::DetKernelType;
}
template<int DIM, int SDIM, int D1D, int Q1D>
DetKernel QuadratureInterpolator::DetKernels::Kernel()
{
if (DIM == 1) { return internal::quadrature_interpolator::Det1D; }
else if (DIM == 2 && SDIM == 2) { return internal::quadrature_interpolator::Det2D<D1D, Q1D>; }
else if (DIM == 2 && SDIM == 3) { return internal::quadrature_interpolator::Det2DSurface<D1D, Q1D>; }
else if (DIM == 3) { return internal::quadrature_interpolator::Det3D<D1D, Q1D>; }
else { MFEM_ABORT(""); }
}
DetKernel QuadratureInterpolator::DetKernels::Fallback(
int DIM, int SDIM, int D1D, int Q1D)
{
if (DIM == 1) { return internal::quadrature_interpolator::Det1D; }
-304
View File
@@ -1,304 +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_QUADINTERP_DET_HPP
#define MFEM_QUADINTERP_DET_HPP
#include "../quadinterpolator.hpp"
#include "../../general/forall.hpp"
#include "../../linalg/dtensor.hpp"
#include "../../fem/kernels.hpp"
#include "../../linalg/kernels.hpp"
namespace mfem
{
namespace internal
{
namespace quadrature_interpolator
{
inline void Det1D(const int NE,
const real_t *b,
const real_t *g,
const real_t *x,
real_t *y,
const int d1d,
const int q1d,
Vector *d_buff = nullptr)
{
MFEM_CONTRACT_VAR(b);
MFEM_CONTRACT_VAR(d_buff);
const auto G = Reshape(g, q1d, d1d);
const auto X = Reshape(x, d1d, NE);
auto Y = Reshape(y, q1d, NE);
mfem::forall(NE, [=] MFEM_HOST_DEVICE (int e)
{
for (int q = 0; q < q1d; q++)
{
real_t u = 0.0;
for (int d = 0; d < d1d; d++)
{
u += G(q, d) * X(d, e);
}
Y(q, e) = u;
}
});
}
template<int T_D1D = 0, int T_Q1D = 0>
inline void Det2D(const int NE,
const real_t *b,
const real_t *g,
const real_t *x,
real_t *y,
const int d1d = 0,
const int q1d = 0,
Vector *d_buff = nullptr)
{
MFEM_CONTRACT_VAR(d_buff);
static constexpr int SDIM = 2;
static constexpr int NBZ = 1;
const int D1D = T_D1D ? T_D1D : d1d;
const int Q1D = T_Q1D ? T_Q1D : q1d;
const auto B = Reshape(b, Q1D, D1D);
const auto G = Reshape(g, Q1D, D1D);
const auto X = Reshape(x, D1D, D1D, SDIM, NE);
auto Y = Reshape(y, Q1D, Q1D, NE);
mfem::forall_2D_batch(NE, Q1D, Q1D, NBZ, [=] MFEM_HOST_DEVICE (int e)
{
constexpr int MQ1 = T_Q1D ? T_Q1D : DofQuadLimits::MAX_Q1D;
constexpr int MD1 = T_D1D ? T_D1D : DofQuadLimits::MAX_D1D;
const int D1D = T_D1D ? T_D1D : d1d;
const int Q1D = T_Q1D ? T_Q1D : q1d;
MFEM_SHARED real_t BG[2][MQ1*MD1];
MFEM_SHARED real_t XY[SDIM][NBZ][MD1*MD1];
MFEM_SHARED real_t DQ[2*SDIM][NBZ][MD1*MQ1];
MFEM_SHARED real_t QQ[2*SDIM][NBZ][MQ1*MQ1];
kernels::internal::LoadX<MD1,NBZ>(e,D1D,X,XY);
kernels::internal::LoadBG<MD1,MQ1>(D1D,Q1D,B,G,BG);
kernels::internal::GradX<MD1,MQ1,NBZ>(D1D,Q1D,BG,XY,DQ);
kernels::internal::GradY<MD1,MQ1,NBZ>(D1D,Q1D,BG,DQ,QQ);
MFEM_FOREACH_THREAD(qy,y,Q1D)
{
MFEM_FOREACH_THREAD(qx,x,Q1D)
{
real_t J[4];
kernels::internal::PullGrad<MQ1,NBZ>(Q1D,qx,qy,QQ,J);
Y(qx,qy,e) = kernels::Det<2>(J);
}
}
});
}
template<int T_D1D = 0, int T_Q1D = 0>
inline void Det2DSurface(const int NE,
const real_t *b,
const real_t *g,
const real_t *x,
real_t *y,
const int d1d = 0,
const int q1d = 0,
Vector *d_buff = nullptr)
{
MFEM_CONTRACT_VAR(d_buff);
static constexpr int SDIM = 3;
static constexpr int NBZ = 1;
const int D1D = T_D1D ? T_D1D : d1d;
const int Q1D = T_Q1D ? T_Q1D : q1d;
const auto B = Reshape(b, Q1D, D1D);
const auto G = Reshape(g, Q1D, D1D);
const auto X = Reshape(x, D1D, D1D, SDIM, NE);
auto Y = Reshape(y, Q1D, Q1D, NE);
mfem::forall_2D_batch(NE, Q1D, Q1D, NBZ, [=] MFEM_HOST_DEVICE (int e)
{
constexpr int MQ1 = T_Q1D ? T_Q1D : DofQuadLimits::MAX_Q1D;
constexpr int MD1 = T_D1D ? T_D1D : DofQuadLimits::MAX_D1D;
const int D1D = T_D1D ? T_D1D : d1d;
const int Q1D = T_Q1D ? T_Q1D : q1d;
const int tidz = MFEM_THREAD_ID(z);
MFEM_SHARED real_t BG[2][MQ1*MD1];
MFEM_SHARED real_t XYZ[SDIM][NBZ][MD1*MD1];
MFEM_SHARED real_t DQ[2*SDIM][NBZ][MD1*MQ1];
kernels::internal::LoadBG<MD1,MQ1>(D1D,Q1D,B,G,BG);
// Load XYZ components
MFEM_FOREACH_THREAD(dy,y,D1D)
{
MFEM_FOREACH_THREAD(dx,x,D1D)
{
for (int d = 0; d < SDIM; ++d)
{
XYZ[d][tidz][dx + dy*D1D] = X(dx,dy,d,e);
}
}
}
MFEM_SYNC_THREAD;
ConstDeviceMatrix B_mat(BG[0], D1D, Q1D);
ConstDeviceMatrix G_mat(BG[1], D1D, Q1D);
// x contraction
MFEM_FOREACH_THREAD(dy,y,D1D)
{
MFEM_FOREACH_THREAD(qx,x,Q1D)
{
for (int d = 0; d < SDIM; ++d)
{
real_t u = 0.0;
real_t v = 0.0;
for (int dx = 0; dx < D1D; ++dx)
{
const real_t xval = XYZ[d][tidz][dx + dy*D1D];
u += xval * G_mat(dx,qx);
v += xval * B_mat(dx,qx);
}
DQ[d][tidz][dy + qx*D1D] = u;
DQ[3 + d][tidz][dy + qx*D1D] = v;
}
}
}
MFEM_SYNC_THREAD;
// y contraction and determinant computation
MFEM_FOREACH_THREAD(qy,y,Q1D)
{
MFEM_FOREACH_THREAD(qx,x,Q1D)
{
real_t J_[6] = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0};
for (int d = 0; d < SDIM; ++d)
{
for (int dy = 0; dy < D1D; ++dy)
{
J_[d] += DQ[d][tidz][dy + qx*D1D] * B_mat(dy,qy);
J_[3 + d] += DQ[3 + d][tidz][dy + qx*D1D] * G_mat(dy,qy);
}
}
DeviceTensor<2> J(J_, 3, 2);
const real_t E = J(0,0)*J(0,0) + J(1,0)*J(1,0) + J(2,0)*J(2,0);
const real_t F = J(0,0)*J(0,1) + J(1,0)*J(1,1) + J(2,0)*J(2,1);
const real_t G = J(0,1)*J(0,1) + J(1,1)*J(1,1) + J(2,1)*J(2,1);
Y(qx,qy,e) = std::sqrt(E*G - F*F);
}
}
});
}
template<int T_D1D = 0, int T_Q1D = 0, bool SMEM = true>
inline void Det3D(const int NE,
const real_t *b,
const real_t *g,
const real_t *x,
real_t *y,
const int d1d = 0,
const int q1d = 0,
Vector *d_buff = nullptr) // used only with SMEM = false
{
constexpr int DIM = 3;
static constexpr int GRID = SMEM ? 0 : 128;
const int D1D = T_D1D ? T_D1D : d1d;
const int Q1D = T_Q1D ? T_Q1D : q1d;
const auto B = Reshape(b, Q1D, D1D);
const auto G = Reshape(g, Q1D, D1D);
const auto X = Reshape(x, D1D, D1D, D1D, DIM, NE);
auto Y = Reshape(y, Q1D, Q1D, Q1D, NE);
real_t *GM = nullptr;
if (!SMEM)
{
const DeviceDofQuadLimits &limits = DeviceDofQuadLimits::Get();
const int max_q1d = T_Q1D ? T_Q1D : limits.MAX_Q1D;
const int max_d1d = T_D1D ? T_D1D : limits.MAX_D1D;
const int max_qd = std::max(max_q1d, max_d1d);
const int mem_size = max_qd * max_qd * max_qd * 9;
d_buff->SetSize(2*mem_size*GRID);
GM = d_buff->Write();
}
mfem::forall_3D_grid(NE, Q1D, Q1D, Q1D, GRID, [=] MFEM_HOST_DEVICE (int e)
{
static constexpr int MQ1 = T_Q1D ? T_Q1D :
(SMEM ? DofQuadLimits::MAX_DET_1D : DofQuadLimits::MAX_Q1D);
static constexpr int MD1 = T_D1D ? T_D1D :
(SMEM ? DofQuadLimits::MAX_DET_1D : DofQuadLimits::MAX_D1D);
static constexpr int MDQ = MQ1 > MD1 ? MQ1 : MD1;
static constexpr int MSZ = MDQ * MDQ * MDQ * 9;
const int bid = MFEM_BLOCK_ID(x);
MFEM_SHARED real_t BG[2][MQ1*MD1];
MFEM_SHARED real_t SM0[SMEM?MSZ:1];
MFEM_SHARED real_t SM1[SMEM?MSZ:1];
real_t *lm0 = SMEM ? SM0 : GM + MSZ*bid;
real_t *lm1 = SMEM ? SM1 : GM + MSZ*(GRID+bid);
real_t (*DDD)[MD1*MD1*MD1] = (real_t (*)[MD1*MD1*MD1]) (lm0);
real_t (*DDQ)[MD1*MD1*MQ1] = (real_t (*)[MD1*MD1*MQ1]) (lm1);
real_t (*DQQ)[MD1*MQ1*MQ1] = (real_t (*)[MD1*MQ1*MQ1]) (lm0);
real_t (*QQQ)[MQ1*MQ1*MQ1] = (real_t (*)[MQ1*MQ1*MQ1]) (lm1);
kernels::internal::LoadX<MD1>(e,D1D,X,DDD);
kernels::internal::LoadBG<MD1,MQ1>(D1D,Q1D,B,G,BG);
kernels::internal::GradX<MD1,MQ1>(D1D,Q1D,BG,DDD,DDQ);
kernels::internal::GradY<MD1,MQ1>(D1D,Q1D,BG,DDQ,DQQ);
kernels::internal::GradZ<MD1,MQ1>(D1D,Q1D,BG,DQQ,QQQ);
MFEM_FOREACH_THREAD(qz,z,Q1D)
{
MFEM_FOREACH_THREAD(qy,y,Q1D)
{
MFEM_FOREACH_THREAD(qx,x,Q1D)
{
real_t J[9];
kernels::internal::PullGrad<MQ1>(Q1D, qx,qy,qz, QQQ, J);
Y(qx,qy,qz,e) = kernels::Det<3>(J);
}
}
}
});
}
} // namespace quadrature_interpolator
} // namespace internal
/// @cond Suppress_Doxygen_warnings
template<int DIM, int SDIM, int D1D, int Q1D>
QuadratureInterpolator::DetKernelType
QuadratureInterpolator::DetKernels::Kernel()
{
if (DIM == 1) { return internal::quadrature_interpolator::Det1D; }
else if (DIM == 2 && SDIM == 2) { return internal::quadrature_interpolator::Det2D<D1D, Q1D>; }
else if (DIM == 2 && SDIM == 3) { return internal::quadrature_interpolator::Det2DSurface<D1D, Q1D>; }
else if (DIM == 3) { return internal::quadrature_interpolator::Det3D<D1D, Q1D>; }
else { MFEM_ABORT(""); }
}
/// @endcond
} // namespace mfem
#endif // MFEM_QUADINTERP_DET_HPP
+8 -25
View File
@@ -96,33 +96,17 @@ void QuadratureSpaceBase::Integrate(VectorCoefficient &coeff,
void QuadratureSpace::ConstructOffsets()
{
MFEM_PERF_FUNCTION;
const int num_elem = mesh.GetNE();
ne = num_elem;
if (mesh.GetNumGeometries(mesh.Dimension()) == 1)
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;
}
void QuadratureSpace::Construct()
@@ -204,7 +188,6 @@ void FaceQuadratureSpace::ConstructOffsets()
{
face_indices.SetSize(num_faces);
offsets.SetSize(num_faces + 1);
ne = num_faces;
int offset = 0;
int f_idx = 0;
for (int i = 0; i < mesh.GetNumFacesWithGhost(); i++)
+4 -26
View File
@@ -30,17 +30,13 @@ protected:
Mesh &mesh; ///< The underlying mesh.
int order; ///< The order of integration rule.
int size; ///< Total number of quadrature points.
int ne; ///< Actual number of entities
mutable Vector weights; ///< Integration weights.
mutable long nodes_sequence = 0; ///< Nodes counter for cache invalidation.
/// @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.
/// @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];
@@ -63,24 +59,6 @@ protected:
void ConstructWeights() const;
public:
/// @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
{
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 { return offsets; }
/// Return the total number of quadrature points.
int GetSize() const { return size; }
@@ -88,7 +66,7 @@ public:
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; }
-1
View File
@@ -503,7 +503,6 @@ 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();
+1 -6
View File
@@ -25,7 +25,7 @@ namespace mfem
ElementRestriction::ElementRestriction(const FiniteElementSpace &f,
ElementDofOrdering e_ordering)
: fes((MFEM_PERF_BEGIN(_MFEM_FUNC_NAME), f)),
: fes(f),
ne(fes.GetNE()),
vdim(fes.GetVDim()),
byvdim(fes.GetOrdering() == Ordering::byVDIM),
@@ -104,13 +104,10 @@ ElementRestriction::ElementRestriction(const FiniteElementSpace &f,
offsets[i] = offsets[i - 1];
}
offsets[0] = 0;
MFEM_PERF_END(_MFEM_FUNC_NAME);
}
void ElementRestriction::Mult(const Vector& x, Vector& y) const
{
MFEM_PERF_FUNCTION;
// Assumes all elements have the same number of dofs
const int nd = dof;
const int vd = vdim;
@@ -155,8 +152,6 @@ void ElementRestriction::AbsMult(const Vector& x, Vector& y) const
template <bool ADD>
void ElementRestriction::TAddMultTranspose(const Vector& x, Vector& y) const
{
MFEM_PERF_FUNCTION;
// Assumes all elements have the same number of dofs
const int nd = dof;
const int vd = vdim;
+27 -9
View File
@@ -5122,32 +5122,33 @@ real_t TMOP_Integrator::GetSurfaceFittingWeight()
void TMOP_Integrator::EnableNormalization(const GridFunction &x)
{
ComputeNormalizationEnergies(x, metric_normal, lim_normal);
ComputeNormalizationEnergies(x, metric_normal, lim_normal, surf_fit_normal);
metric_normal = 1.0 / metric_normal;
lim_normal = 1.0 / lim_normal;
//if (surf_fit_gf) { surf_fit_normal = 1.0 / surf_fit_normal; }
if (surf_fit_gf || surf_fit_pos) { surf_fit_normal = lim_normal; }
}
#ifdef MFEM_USE_MPI
void TMOP_Integrator::ParEnableNormalization(const ParGridFunction &x)
{
real_t loc[2];
ComputeNormalizationEnergies(x, loc[0], loc[1]);
real_t rdc[2];
MPI_Allreduce(loc, rdc, 2, MPITypeMap<real_t>::mpi_type, MPI_SUM,
real_t loc[3];
ComputeNormalizationEnergies(x, loc[0], loc[1], loc[2]);
real_t rdc[3];
MPI_Allreduce(loc, rdc, 3, MPITypeMap<real_t>::mpi_type, MPI_SUM,
x.ParFESpace()->GetComm());
metric_normal = 1.0 / rdc[0];
lim_normal = 1.0 / rdc[1];
// if (surf_fit_gf) { surf_fit_normal = 1.0 / rdc[2]; }
if (surf_fit_gf || surf_fit_pos) { surf_fit_normal = lim_normal; }
}
#endif
void TMOP_Integrator::ComputeNormalizationEnergies(const GridFunction &x,
real_t &metric_energy,
real_t &lim_energy)
real_t &lim_energy,
real_t &surf_fit_gf_energy)
{
metric_energy = 0.0;
lim_energy = 0.0;
if (PA.enabled)
{
MFEM_VERIFY(PA.E.Size() > 0, "Must be called after AssemblePA!");
@@ -5190,6 +5191,9 @@ void TMOP_Integrator::ComputeNormalizationEnergies(const GridFunction &x,
Jpr.SetSize(dim);
Jpt.SetSize(dim);
metric_energy = 0.0;
lim_energy = 0.0;
surf_fit_gf_energy = 0.0;
for (int i = 0; i < fes->GetNE(); i++)
{
const FiniteElement *fe = fes->GetFE(i);
@@ -5221,7 +5225,21 @@ void TMOP_Integrator::ComputeNormalizationEnergies(const GridFunction &x,
lim_energy += weight;
}
// TODO: Normalization of the surface fitting term.
// Normalization of the surface fitting term.
if (surf_fit_gf)
{
Array<int> dofs;
Vector sigma_e;
surf_fit_gf->FESpace()->GetElementDofs(i, dofs);
surf_fit_gf->GetSubVector(dofs, sigma_e);
for (int s = 0; s < dofs.Size(); s++)
{
if ((*surf_fit_marker)[dofs[s]] == true)
{
surf_fit_gf_energy += sigma_e(s) * sigma_e(s);
}
}
}
}
// Cases when integration is not over the target element, or when the
+2 -1
View File
@@ -2038,7 +2038,8 @@ protected:
} PA;
void ComputeNormalizationEnergies(const GridFunction &x,
real_t &metric_energy, real_t &lim_energy);
real_t &metric_energy, real_t &lim_energy,
real_t &surf_fit_gf_energy);
void AssembleElementVectorExact(const FiniteElement &el,
ElementTransformation &T,
+17 -196
View File
@@ -13,7 +13,6 @@
#include "bilinearform.hpp"
#include "pbilinearform.hpp"
#include "../general/forall.hpp"
#include "kernels.hpp"
namespace mfem
{
@@ -2323,76 +2322,6 @@ 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)
{
MFEM_PERF_FUNCTION;
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)
@@ -2474,9 +2403,9 @@ void Prolongation3D(const int NE, const int D1D, const int Q1D,
});
}
void ProlongationTranspose2D(const int NE, const int D1D, const int Q1D,
const Vector& localH, Vector& localL,
const Array<real_t>& Bt, const Vector& mask)
void Restriction2D(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);
@@ -2519,80 +2448,9 @@ void ProlongationTranspose2D(const int NE, const int D1D, const int Q1D,
}
});
}
template <int DLO, int DHI>
static void SmemProlongationTranspose3D(
const int NE, const Vector& localH, Vector& localL,
const Array<real_t>& bt, const Vector& mask)
{
MFEM_PERF_FUNCTION;
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)
void Restriction3D(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);
@@ -2660,15 +2518,11 @@ void ProlongationTranspose3D(const int NE, const int D1D, const int Q1D,
}
});
}
} // namespace TransferKernels
void TensorProductPRefinementTransferOperator::Mult(const Vector& x,
Vector& y) const
{
MFEM_PERF_FUNCTION;
using namespace TransferKernels;
if (lFESpace.GetMesh()->GetNE() == 0)
{
return;
@@ -2677,25 +2531,11 @@ void TensorProductPRefinementTransferOperator::Mult(const Vector& x,
elem_restrict_lex_l->Mult(x, localL);
if (dim == 2)
{
Prolongation2D(NE, D1D, Q1D, localL, localH, B, mask);
TransferKernels::Prolongation2D(NE, D1D, Q1D, localL, localH, B, mask);
}
else if (dim == 3)
{
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;
}
TransferKernels::Prolongation3D(NE, D1D, Q1D, localL, localH, B, mask);
}
else
{
@@ -2709,9 +2549,6 @@ void TensorProductPRefinementTransferOperator::Mult(const Vector& x,
void TensorProductPRefinementTransferOperator::MultTranspose(const Vector& x,
Vector& y) const
{
MFEM_PERF_FUNCTION;
using namespace TransferKernels;
if (lFESpace.GetMesh()->GetNE() == 0)
{
return;
@@ -2720,25 +2557,11 @@ void TensorProductPRefinementTransferOperator::MultTranspose(const Vector& x,
elem_restrict_lex_h->Mult(x, localH);
if (dim == 2)
{
ProlongationTranspose2D(NE, D1D, Q1D, localH, localL, Bt, mask);
TransferKernels::Restriction2D(NE, D1D, Q1D, localH, localL, Bt, mask);
}
else if (dim == 3)
{
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;
}
TransferKernels::Restriction3D(NE, D1D, Q1D, localH, localL, Bt, mask);
}
else
{
@@ -2760,20 +2583,20 @@ TrueTransferOperator::TrueTransferOperator(const FiniteElementSpace& lFESpace_,
P = lFESpace.GetProlongationMatrix();
R = hFESpace.IsVariableOrder() ? hFESpace.GetHpRestrictionMatrix() :
hFESpace.GetRestrictionOperator();
hFESpace.GetRestrictionMatrix();
// P and R can be both null
// P can be null and R not null
// If P is not null it is assumed that R is not null as well
if (P) { MFEM_VERIFY(R, "Both P and R have to be not NULL") }
if (!IsIdentityProlongation(P))
if (P)
{
tmpL.SetSize(lFESpace_.GetVSize());
tmpH.SetSize(hFESpace_.GetVSize());
}
// P can be null and R not null
else if (!IsIdentityProlongation(R))
else if (R)
{
tmpH.SetSize(hFESpace_.GetVSize());
}
@@ -2786,14 +2609,13 @@ TrueTransferOperator::~TrueTransferOperator()
void TrueTransferOperator::Mult(const Vector& x, Vector& y) const
{
MFEM_PERF_FUNCTION;
if (!IsIdentityProlongation(P))
if (P)
{
P->Mult(x, tmpL);
localTransferOperator->Mult(tmpL, tmpH);
R->Mult(tmpH, y);
}
else if (!IsIdentityProlongation(R))
else if (R)
{
localTransferOperator->Mult(x, tmpH);
R->Mult(tmpH, y);
@@ -2806,14 +2628,13 @@ void TrueTransferOperator::Mult(const Vector& x, Vector& y) const
void TrueTransferOperator::MultTranspose(const Vector& x, Vector& y) const
{
MFEM_PERF_FUNCTION;
if (!IsIdentityProlongation(P))
if (P)
{
R->MultTranspose(x, tmpH);
localTransferOperator->MultTranspose(tmpH, tmpL);
P->MultTranspose(tmpL, y);
}
else if (!IsIdentityProlongation(R))
else if (R)
{
R->MultTranspose(x, tmpH);
localTransferOperator->MultTranspose(tmpH, y);
+4 -1
View File
@@ -621,6 +621,9 @@ 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;
@@ -639,7 +642,7 @@ private:
const FiniteElementSpace& lFESpace;
const FiniteElementSpace& hFESpace;
const Operator * P = nullptr;
const Operator * R = nullptr;
const SparseMatrix * R = nullptr;
TransferOperator* localTransferOperator;
mutable Vector tmpL;
mutable Vector tmpH;
+9 -83
View File
@@ -14,98 +14,24 @@
#include "../config/config.hpp"
#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 <caliper/cali.h>
#include <caliper/cali-manager.h>
#endif
namespace mfem
{
namespace internal
{
extern int annotation_sync_stream; // defined in globals.cpp
extern int annotation_sync_mpi; // defined in globals.cpp
#ifdef MFEM_USE_CALIPER
inline void AnnotationSync()
{
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
{
::cali::Function cali_func;
FunctionAnnotation(const char *fname)
: cali_func((AnnotationSync(), fname)) { }
~FunctionAnnotation() { AnnotationSync(); }
};
struct ScopeAnnotation
{
::cali::ScopeAnnotation cali_scope;
ScopeAnnotation(const char *name)
: cali_scope((AnnotationSync(), 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_FUNCTION CALI_CXX_MARK_FUNCTION
#define MFEM_PERF_BEGIN(s) CALI_MARK_BEGIN(s)
#define MFEM_PERF_END(s) CALI_MARK_END(s)
#define MFEM_PERF_SCOPE(name) \
mfem::internal::ScopeAnnotation \
MFEM_CONCAT(mfem_scope_annotation_,__LINE__)(name)
cali::Annotation::Guard cali_autogenerated_guard_name(cali::Annotation("function").begin(std::string(name).c_str()))
#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))
#else // #ifdef MFEM_USE_CALIPER
#else
#define MFEM_PERF_FUNCTION
#define MFEM_PERF_BEGIN(s) ((void)(0))
#define MFEM_PERF_BEGIN(s)
#define MFEM_PERF_END(s)
#define MFEM_PERF_SCOPE(name)
#define MFEM_PERF_SYNC_STREAM(b)
#define MFEM_PERF_SYNC_MPI(b)
#define MFEM_PERF_SYNC(b)
#endif
#endif // #ifdef MFEM_USE_CALIPER
#endif // MFEM_ANNOTATION_HPP
#endif
+1 -5
View File
@@ -326,11 +326,7 @@ public:
the Size to match this Capacity after this.*/
template <typename U>
inline void CopyFrom(const U *src)
{
if (!begin() || size == 0) { return; }
MFEM_ASSERT(begin() && src, "Error in Array::CopyFrom");
std::memcpy(begin(), src, MemoryUsage());
}
{ std::memcpy(begin(), src, MemoryUsage()); }
/// STL-like begin. Returns pointer to the first element of the array.
inline T* begin() { return data; }
+6 -7
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,18 +55,17 @@
#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
#define MFEM_THREAD_ID(k) 0
#define MFEM_THREAD_SIZE(k) 1
#define MFEM_FOREACH_THREAD(i,k,N) for(int i=0; i<N; i++)
#define MFEM_FOREACH_THREAD_DIRECT(i,k,N) MFEM_FOREACH_THREAD(i,k,N)
#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 +93,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)
+24 -59
View File
@@ -23,9 +23,13 @@
#include <mpi.h>
#include <cstdint>
// 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
// 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
namespace mfem
{
@@ -421,24 +425,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.
@@ -603,50 +599,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;
@@ -655,6 +611,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 -6
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,23 +37,22 @@
__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
#define MFEM_THREAD_ID(k) threadIdx.k
#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);
+1 -53
View File
@@ -16,7 +16,6 @@
#include "../fem/ceed/interface/util.hpp"
#endif
#ifdef MFEM_USE_MPI
#include "communication.hpp"
#include "../linalg/hypre.hpp"
#endif
@@ -146,27 +145,6 @@ Device::Device()
Configure(device);
device_env = true;
}
if (GetEnv("MFEM_GPU_AWARE_MPI"))
{
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()
@@ -218,29 +196,6 @@ void Device::Configure(const std::string &device, const int device_id)
{
bmap[internal::backend_name[i]] = internal::backend_list[i];
}
// auto-detect GPU configurations
// assumes only one of HIP or CUDA are available
#ifdef MFEM_USE_HIP
bmap["gpu"] = Backend::HIP;
#ifdef MFEM_USE_RAJA
bmap["raja-gpu"] = Backend::RAJA_HIP;
#endif
#ifdef MFEM_USE_CEED
bmap["ceed-gpu"] = Backend::CEED_HIP;
#endif
// no OCCA+HIP?
#elif defined(MFEM_USE_CUDA)
bmap["gpu"] = Backend::CUDA;
#ifdef MFEM_USE_RAJA
bmap["raja-gpu"] = Backend::RAJA_CUDA;
#endif
#ifdef MFEM_USE_CEED
bmap["ceed-gpu"] = Backend::CEED_CUDA;
#endif
#ifdef MFEM_USE_OCCA
bmap["occa-gpu"] = Backend::OCCA_CUDA;
#endif
#endif
std::string device_option;
std::string::size_type beg = 0, end;
while (1)
@@ -358,13 +313,6 @@ void Device::Print(std::ostream &os)
{
os << ',' << MemoryTypeName[static_cast<int>(device_mem_type)];
}
#ifdef MFEM_USE_MPI
if (Allows(Backend::DEVICE_MASK) &&
Mpi::IsInitialized() && !Mpi::IsFinalized())
{
os << "\nUse GPU-aware MPI: " << (GetGPUAwareMPI() ? "yes" : "no");
}
#endif
os << std::endl;
}
@@ -631,7 +579,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 -7
View File
@@ -198,10 +198,6 @@ public:
'ceed-hip', 'hip', 'debug',
'occa-omp', 'raja-omp', 'omp',
'ceed-cpu', 'occa-cpu', 'raja-cpu', 'cpu'.
- The following backend aliases are also available: 'ceed-gpu',
'occa-gpu', 'raja-gpu', and 'gpu' where they alias their respective
'*-cuda' or '*-hip' backends depending on the MFEM build-time
configuration.
- Multiple backends can be configured at the same time.
- Only one 'occa-*' backend can be configured at a time.
- The backend 'occa-cuda' enables the 'cuda' backend unless 'raja-cuda'
@@ -297,9 +293,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 // MFEM_ERROR_HPP
#endif
+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)
{
-3
View File
@@ -31,9 +31,6 @@ 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
}
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();
+5 -7
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,20 +37,18 @@
__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
#define MFEM_THREAD_ID(k) hipThreadIdx_ ##k
#define MFEM_THREAD_SIZE(k) hipBlockDim_ ##k
#define MFEM_FOREACH_THREAD(i,k,N) \
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__)
for(int i=hipThreadIdx_ ##k; i<N; i+=hipBlockDim_ ##k)
#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;
}
+5 -9
View File
@@ -657,8 +657,7 @@ 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,
@@ -742,7 +741,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,
@@ -897,7 +896,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;
@@ -982,7 +980,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_((void*)h_ptr);
MemoryType h_ptr_mt = MemoryManager::GetHostMemoryType_(h_ptr);
MFEM_VERIFY(h_mt == h_ptr_mt,
"h_mt = " << (int)h_mt << ", h_ptr_mt = " << (int)h_ptr_mt);
}
@@ -990,8 +988,7 @@ 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_((void*)ptr, (void*)ptr, bytes, h_mt, own, false,
flags);
MemoryManager::Register_(ptr, ptr, bytes, h_mt, own, false, flags);
}
}
@@ -1060,8 +1057,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
)
{

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