Skip to content

Commit

Permalink
Merge github.com:xcompact3d/x3d2 into 2decomp-integration
Browse files Browse the repository at this point in the history
  • Loading branch information
Nanoseb committed Oct 17, 2024
2 parents eb3883b + b92d919 commit 9b9bc35
Show file tree
Hide file tree
Showing 31 changed files with 1,340 additions and 441 deletions.
128 changes: 128 additions & 0 deletions .gitlab-ci.yml
Original file line number Diff line number Diff line change
@@ -0,0 +1,128 @@
workflow:
rules:
# Run pipeline on tags for the main project
- if: $CI_COMMIT_TAG && $CI_PROJECT_PATH == "CFD-Xing/x3d2"
# Run pipeline on the default branch for the main project
- if: $CI_COMMIT_BRANCH == "main" && $CI_PROJECT_PATH == "CFD-Xing/x3d2"
# DO NOT run pipeline if Draft
- if: $CI_PIPELINE_SOURCE == "external_pull_request_event" && $CI_MERGE_REQUEST_TITLE =~ /^(\[Draft\]|\(Draft\)|Draft:).*/
when: never
# Other merge requests trigger pipelines
- if: $CI_PIPELINE_SOURCE == "external_pull_request_event"
changes:
compare_to: 'refs/heads/main'
paths:
- "src/**/*"
- "examples/**/*"
- "tests/**/*"
- .gitlab-ci.yml

stages:
- build-and-test
- check-policies
- build-docs

.build-and-test-template: &build-and-test-template
image: ubuntu:22.04
stage: build-and-test
timeout: 1h
script:
- echo "Setup environment"
- apt update
- apt install -y environment-modules
- echo "/apps/modules" >> /etc/environment-modules/modulespath
- apt-get update
- apt-get install -y ccache
- apt-get install -y cmake
- apt-get install -y gcc
- apt-get install -y infiniband-diags ibverbs-utils
- apt-get install -y libibverbs-dev libfabric1 libfabric-dev libpsm2-dev
- apt-get install -y openmpi-bin openmpi-common libopenmpi-dev libgtk2.0-dev
- apt-get install -y librdmacm-dev libpsm2-dev
- . /etc/profile.d/modules.sh
- ccache -s && ccache -M 5G
- echo "Display GCC version number"
- gcc --version
- echo "Display nvcc version number"
- module avail
#- module load cuda
#- nvcc --version
- echo "Configure Debug build"
- FC=mpif90 cmake -S . -B build -DCMAKE_BUILD_TYPE=Debug
- echo "Build with Debug (strict) flags"
- make -C build
- echo "Configure Release build"
- FC=mpif90 cmake -S . -B build -DCMAKE_BUILD_TYPE=Release
- echo "Build tests"
- make -C build
- echo "Run the tests"
- export OMPI_ALLOW_RUN_AS_ROOT=1
- export OMPI_ALLOW_RUN_AS_ROOT_CONFIRM=1
- make -C build test
tags:
- gpu

.check-formatting-template: &check-formatting-template
image: python:3.9.19-bullseye
stage: check-policies
variables:
FPRETTIFY_COMMAND: fprettify --config .fprettify.ini --diff --recursive src
before_script:
- pip install fprettify
script:
- cd $CI_PROJECT_DIR
- $FPRETTIFY_COMMAND &> fprettify.log
- if [[ ! -z "$(cat fprettify.log)" ]]; then echo "::warning::Code formatting issues detected. See log for details."; exit 1; fi
allow_failure: false
timeout: 15m
artifacts:
expire_in: 1 month
when: on_failure
paths:
- fprettify.log

.build-docs-template: &build-docs-template
image: python:3.9.19-slim-bookworm
stage: build-docs
variables:
PUBLISH_DIR: api/
FORD_OUTPUT_DIR: api-docs
FORD_CFG: ford.md
script:
- cd $CI_PROJECT_DIR
- apt update
- apt install -y git
- echo Install sphinx, ford and ghp-import
- pip install -r docs/docs-requirements.txt
- echo Install graphviz
- apt install graphviz -y
- echo Build sphinx docs
- apt-get install make -y
- make -C docs html
- echo Deploy sphinx docs
- git config user.name 'github-action'
- git config user.email 'github-action'
- git remote add upstream git@github.com:xcompact3d/x3d2.git
- ghp-import -m 'Update sphinx docs' --push --remote upstream --branch gh-pages docs/build/html --no-jekyll --force
- echo Build API docs with ford
- ford $FORD_CFG -o $FORD_OUTPUT_DIR
- echo Deploy api-docs
- ghp-import -m 'Update API docs' --prefix $PUBLISH_DIR --push --remote upstream --branch gh-pages $FORD_OUTPUT_DIR --no-jekyll --force
rules:
- if: '$CI_COMMIT_BRANCH == "main"'
when: on_success
- when: never
allow_failure: false
timeout: 15m

build-and-test:
<<: *build-and-test-template
needs: []

check-formatting:
<<: *check-formatting-template
needs: []

#build-docs:
# <<: *build-docs-template
# needs: []
28 changes: 28 additions & 0 deletions examples/TGV/input.x3d
Original file line number Diff line number Diff line change
@@ -0,0 +1,28 @@
&domain_params
! Global number of cells in each direction
L_global = 6.283185307179586d0, 6.283185307179586d0, 6.283185307179586d0

! Global domain dimensions
dims_global = 256, 256, 256

! Domain decomposition in each direction
nproc_dir = 1, 1, 2

! BC options are 'periodic' | 'neumann' | 'dirichlet'
BC_x = 'periodic', 'periodic'
BC_y = 'periodic', 'periodic'
BC_z = 'periodic', 'periodic'
/End

&solver_params
Re = 1600d0
time_intg = 'AB3' ! 'AB[1-4]' | 'RK[1-4]'
dt = 0.001d0
n_iters = 20000
n_output = 100
poisson_solver_type = 'FFT' ! 'FFT' | 'CG'
der1st_scheme = 'compact6'
der2nd_scheme = 'compact6' ! 'compact6' | 'compact6-hyperviscous'
interpl_scheme = 'classic' ! 'classic' | 'optimised' | 'aggressive'
stagder_scheme = 'compact6'
/End
6 changes: 6 additions & 0 deletions src/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -2,6 +2,9 @@ set(SRC
allocator.f90
backend.f90
common.f90
field.f90
mesh.f90
ordering.f90
poisson_fft.f90
solver.f90
tdsops.f90
Expand All @@ -11,6 +14,7 @@ set(SRC
decomp.f90
field.f90
par_grid.f90
vector_calculus.f90
omp/backend.f90
omp/common.f90
omp/kernels/distributed.f90
Expand All @@ -23,9 +27,11 @@ set(CUDASRC
cuda/backend.f90
cuda/common.f90
cuda/exec_dist.f90
cuda/exec_thom.f90
cuda/kernels/distributed.f90
cuda/kernels/reorder.f90
cuda/kernels/spectral_processing.f90
cuda/kernels/thomas.f90
cuda/poisson_fft.f90
cuda/sendrecv.f90
cuda/tdsops.f90
Expand Down
15 changes: 6 additions & 9 deletions src/backend.f90
Original file line number Diff line number Diff line change
Expand Up @@ -27,7 +27,6 @@ module m_base_backend
real(dp) :: nu
class(mesh_t), pointer :: mesh
class(allocator_t), pointer :: allocator
class(dirps_t), pointer :: xdirps, ydirps, zdirps
class(poisson_fft_t), pointer :: poisson_fft
contains
procedure(transeq_ders), deferred :: transeq_x
Expand Down Expand Up @@ -68,22 +67,20 @@ end subroutine transeq_ders
end interface

abstract interface
subroutine tds_solve(self, du, u, dirps, tdsops)
!! transeq equation obtains the derivatives direction by
!! direction, and the exact algorithm used to obtain these
!! derivatives are decided at runtime. Backend implementations
!! are responsible from directing calls to transeq_ders into
!! the correct algorithm.
subroutine tds_solve(self, du, u, tdsops)
!! transeq equation obtains the derivatives direction by
!! direction, and the exact algorithm used to obtain these
!! derivatives are decided at runtime. Backend implementations
!! are responsible from directing calls to tds_solve to the
!! correct algorithm.
import :: base_backend_t
import :: field_t
import :: dirps_t
import :: tdsops_t
implicit none

class(base_backend_t) :: self
class(field_t), intent(inout) :: du
class(field_t), intent(in) :: u
type(dirps_t), intent(in) :: dirps
class(tdsops_t), intent(in) :: tdsops
end subroutine tds_solve
end interface
Expand Down
26 changes: 10 additions & 16 deletions src/common.f90
Original file line number Diff line number Diff line change
Expand Up @@ -10,28 +10,22 @@ module m_common
RDR_X2C = 14, RDR_Y2C = 24, RDR_Z2C = 34
integer, parameter :: DIR_X = 1, DIR_Y = 2, DIR_Z = 3, DIR_C = 4
integer, parameter :: POISSON_SOLVER_FFT = 0, POISSON_SOLVER_CG = 1
integer, parameter :: VERT = 1, & ! Vertex centered data
CELL = 2, & ! Cell centered data
X_FACE = 11, & ! Data on faces normal to X
Y_FACE = 12, & ! Data on faces normal to Y
Z_FACE = 13, & ! Data on faces normal to Z
X_EDGE = 101, & ! Data on edges along X
Y_EDGE = 102, & ! Data on edges along Y
Z_EDGE = 103, & ! Data on edges along Z
none = -1 ! The location of data isn't specified
integer, parameter :: VERT = 0000, & ! Vertex centered data
CELL = 1110, & ! Cell centered data
X_FACE = 0110, & ! Data on faces normal to X
Y_FACE = 1010, & ! Data on faces normal to Y
Z_FACE = 1100, & ! Data on faces normal to Z
X_EDGE = 1000, & ! Data on edges along X
Y_EDGE = 0100, & ! Data on edges along Y
Z_EDGE = 0010, & ! Data on edges along Z
none = -0001 ! The location of data isn't specified
integer, parameter :: BC_PERIODIC = 0, BC_NEUMANN = 1, BC_DIRICHLET = 2
integer, protected :: &
rdr_map(4, 4) = reshape([0, RDR_Y2X, RDR_Z2X, RDR_C2X, &
RDR_X2Y, 0, RDR_Z2Y, RDR_C2Y, &
RDR_X2Z, RDR_Y2Z, 0, RDR_C2Z, &
RDR_X2C, RDR_Y2C, RDR_Z2C, 0], shape=[4, 4])

type :: globs_t
real(dp) :: nu, dt
integer :: n_iters, n_output
character(len=20) :: BC_x_s, BC_x_e, BC_y_s, BC_y_e, BC_z_s, BC_z_e
integer :: poisson_solver_type
end type globs_t

contains

pure subroutine get_dirs_from_rdr(dir_from, dir_to, rdr_dir)
Expand Down
14 changes: 6 additions & 8 deletions src/cuda/backend.f90
Original file line number Diff line number Diff line change
Expand Up @@ -5,7 +5,7 @@ module m_cuda_backend

use m_allocator, only: allocator_t, field_t
use m_base_backend, only: base_backend_t
use m_common, only: dp, globs_t, &
use m_common, only: dp, &
RDR_X2Y, RDR_X2Z, RDR_Y2X, RDR_Y2Z, RDR_Z2X, RDR_Z2Y, &
RDR_C2X, RDR_C2Y, RDR_C2Z, RDR_X2C, RDR_Y2C, RDR_Z2C, &
DIR_X, DIR_Y, DIR_Z, DIR_C, VERT
Expand Down Expand Up @@ -352,35 +352,33 @@ subroutine transeq_cuda_thom(self, du, dv, dw, u, v, w, dirps)

end subroutine transeq_cuda_thom

subroutine tds_solve_cuda(self, du, u, dirps, tdsops)
subroutine tds_solve_cuda(self, du, u, tdsops)
implicit none

class(cuda_backend_t) :: self
class(field_t), intent(inout) :: du
class(field_t), intent(in) :: u
type(dirps_t), intent(in) :: dirps
class(tdsops_t), intent(in) :: tdsops

type(dim3) :: blocks, threads

! Check if direction matches for both in/out fields and dirps
if (dirps%dir /= du%dir .or. u%dir /= du%dir) then
error stop 'DIR mismatch between fields and dirps in tds_solve.'
if (u%dir /= du%dir) then
error stop 'DIR mismatch between fields in tds_solve.'
end if

blocks = dim3(self%mesh%get_n_groups(u), 1, 1); threads = dim3(SZ, 1, 1)

call tds_solve_dist(self, du, u, dirps, tdsops, blocks, threads)
call tds_solve_dist(self, du, u, tdsops, blocks, threads)

end subroutine tds_solve_cuda

subroutine tds_solve_dist(self, du, u, dirps, tdsops, blocks, threads)
subroutine tds_solve_dist(self, du, u, tdsops, blocks, threads)
implicit none

class(cuda_backend_t) :: self
class(field_t), intent(inout) :: du
class(field_t), intent(in) :: u
type(dirps_t), intent(in) :: dirps
class(tdsops_t), intent(in) :: tdsops
type(dim3), intent(in) :: blocks, threads

Expand Down
37 changes: 37 additions & 0 deletions src/cuda/exec_thom.f90
Original file line number Diff line number Diff line change
@@ -0,0 +1,37 @@
module m_cuda_exec_thom
use cudafor

use m_common, only: dp
use m_cuda_kernels_thom, only: der_univ_thom, der_univ_thom_per
use m_cuda_tdsops, only: cuda_tdsops_t

implicit none

contains

subroutine exec_thom_tds_compact(du, u, tdsops, blocks, threads)
implicit none

real(dp), device, dimension(:, :, :), intent(out) :: du
real(dp), device, dimension(:, :, :), intent(in) :: u
type(cuda_tdsops_t), intent(in) :: tdsops
type(dim3), intent(in) :: blocks, threads

if (tdsops%periodic) then
call der_univ_thom_per<<<blocks, threads>>>( & !&
du, u, tdsops%coeffs_dev, tdsops%tds_n, tdsops%alpha, &
tdsops%thom_f_dev, tdsops%thom_s_dev, &
tdsops%thom_w_dev, tdsops%thom_p_dev &
)
else
call der_univ_thom<<<blocks, threads>>>( & !&
du, u, &
tdsops%coeffs_s_dev, tdsops%coeffs_e_dev, tdsops%coeffs_dev, &
tdsops%tds_n, tdsops%thom_f_dev, tdsops%thom_s_dev, &
tdsops%thom_w_dev &
)
end if

end subroutine exec_thom_tds_compact

end module m_cuda_exec_thom
Loading

0 comments on commit 9b9bc35

Please sign in to comment.