Compare commits

..

12 Commits

Author SHA1 Message Date
d11eaa2242 Optimize bssn_rhs.f90: Fuse loops for metric inversion and Christoffel symbols to improve cache locality 2026-01-21 11:22:33 +08:00
ef96766e22 优化 compute_rhs_bssn 热点路径并加入 NaN 检查开关
- 用 DEBUG_NAN_CHECK 宏按需启用 NaN 检查,并在输入/宏生成器中新增 Debug_NaN_Check 配置
  - 逆度量改为先求行列式再乘法展开,减少除法;并在 Gam^i/Christoffel 处提取公共子表达式
  - 预置批量 fderivs 辅助例程,便于后续矢量化/合并导数计算
  - 将默认 MPI_processes 调整为 8

  变更涉及:

  - AMSS_NCKU_source/bssn_rhs.f90
  - generate_macrodef.py
  - AMSS_NCKU_Input.py
  - AMSS_NCKU_Input_Mini.py
  - inputfile_example/AMSS_NCKU_Input.py
  - AMSS_NCKU_source/diff_new.f90

TODO: fmisc.f90 polint()
2026-01-20 19:37:26 +08:00
ae7b77e44c Setup GW150914-mini test case for laptop development
- Add AMSS_NCKU_Input_Mini.py with reduced grid resolution and MPI processes
- Add AMSS_NCKU_MiniProgram.py launcher with automatic configuration swapping
- Update makefile_and_run.py to reduce build jobs and remove CPU binding for laptop
- Update .gitignore to exclude GW150914-mini output directory
2026-01-20 00:31:40 +08:00
26c81d8e81 makefile updated 2026-01-19 23:53:16 +08:00
CGH0S7
9deeda9831 Refactor verification method and optimize numerical kernels with oneMKL BLAS
This commit transitions the verification approach from post-Newtonian theory
   comparison to regression testing against baseline simulations, and optimizes
   critical numerical kernels using Intel oneMKL BLAS routines.

   Verification Changes:
   - Replace PN theory-based RMS calculation with trajectory-based comparison
   - Compare optimized results against baseline (GW150914-origin) on XY plane
   - Compute RMS independently for BH1 and BH2, report maximum as final metric
   - Update documentation to reflect new regression test methodology

   Performance Optimizations:
   - Replace manual vector operations with oneMKL BLAS routines:
     * norm2() and scalarproduct() now use cblas_dnrm2/cblas_ddot (C++)
     * L2 norm calculations use DDOT for dot products (Fortran)
     * Interpolation weighted sums use DDOT (Fortran)
   - Disable OpenMP threading (switch to sequential MKL) for better performance

   Build Configuration:
   - Switch from lmkl_intel_thread to lmkl_sequential
   - Remove -qopenmp flags from compiler options
   - Maintain aggressive optimization flags (-O3, -xHost, -fp-model fast=2, -fma)

   Other Changes:
   - Update .gitignore for GW150914-origin, docs, and temporary files
2026-01-18 14:25:21 +08:00
CGH0S7
3a7bce3af2 Update Intel oneAPI configuration and CPU binding settings
- Update makefile.inc with Intel oneAPI compiler flags and oneMKL linking
   - Configure taskset CPU binding to use nohz_full cores (4-55, 60-111)
   - Set build parallelism to 104 jobs for faster compilation
   - Update MPI process count to 48 in input configuration
2026-01-17 20:41:02 +08:00
CGH0S7
c6945bb095 Rename verify_accuracy.py to AMSS_NCKU_Verify_ASC26.py and improve visual output 2026-01-17 14:54:33 +08:00
CGH0S7
0d24f1503c Add accuracy verification script for GW150914 simulation
- Verify RMS error < 1% (black hole trajectory vs. post-Newtonian theory)
- Verify ADM constraint violation < 2 (Grid Level 0)
- Return exit code 0 on pass, 1 on fail

Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
2026-01-17 00:37:30 +08:00
CGH0S7
cb252f5ea2 Optimize numerical algorithms with Intel oneMKL
- FFT.f90: Replace hand-written Cooley-Tukey FFT with oneMKL DFTI
   - ilucg.f90: Replace manual dot product loop with BLAS DDOT
   - gaussj.C: Replace Gauss-Jordan elimination with LAPACK dgesv/dgetri
   - makefile.inc: Add MKL include paths and library linking

   All optimizations maintain mathematical equivalence and numerical precision.
2026-01-16 10:58:11 +08:00
CGH0S7
7a76cbaafd Add numactl CPU binding to avoid cores 0-3 and 56-59
Bind all computation processes (ABE, ABEGPU, TwoPunctureABE) to
   CPU cores 4-55 and 60-111 using numactl --physcpubind to prevent
   interference with system processes on reserved cores.
2026-01-16 10:24:46 +08:00
CGH0S7
57a7376044 Switch compiler toolchain from GCC to Intel oneAPI
- makefile.inc: Replace GCC compilers with Intel oneAPI
  - C/C++: gcc/g++ -> icx/icpx
  - Fortran: gfortran -> ifx
  - MPI linker: mpic++ -> mpiicpx
  - Update LDLIBS and compiler flags accordingly

- macrodef.h: Fix include path (microdef.fh -> macrodef.fh)

Requires: source /home/intel/oneapi/setvars.sh before build
2026-01-15 16:32:12 +08:00
cd5ceaa15f main branch updated 2026-01-14 08:55:53 +08:00
22 changed files with 2563 additions and 1482 deletions

4
.gitignore vendored
View File

@@ -1,3 +1,7 @@
__pycache__ __pycache__
GW150914 GW150914
GW150914-origin
GW150914-mini
docs
*.tmp

View File

@@ -16,12 +16,14 @@ import numpy
File_directory = "GW150914" ## output file directory File_directory = "GW150914" ## output file directory
Output_directory = "binary_output" ## binary data file directory Output_directory = "binary_output" ## binary data file directory
## The file directory name should not be too long ## The file directory name should not be too long
MPI_processes = 96 ## number of mpi processes used in the simulation MPI_processes = 8 ## number of mpi processes used in the simulation
GPU_Calculation = "no" ## Use GPU or not
## (prefer "no" in the current version, because the GPU part may have bugs when integrated in this Python interface)
CPU_Part = 1.0
GPU_Part = 0.0
Debug_NaN_Check = 0 ## enable NaN checks in compute_rhs_bssn: 0 (off) or 1 (on)
GPU_Calculation = "yes" ## Use GPU or not
## GPU support has been updated for CUDA 13
CPU_Part = 0.0
GPU_Part = 1.0
################################################# #################################################

233
AMSS_NCKU_Input_Mini.py Normal file
View File

@@ -0,0 +1,233 @@
#################################################
##
## This file provides the input parameters required for numerical relativity.
## XIAOQU
## 2024/03/19 --- 2025/09/14
## Modified for GW150914-mini test case
##
#################################################
import numpy
#################################################
## Setting MPI processes and the output file directory
File_directory = "GW150914-mini" ## output file directory
Output_directory = "binary_output" ## binary data file directory
## The file directory name should not be too long
MPI_processes = 4 ## number of mpi processes used in the simulation (Reduced for laptop)
GPU_Calculation = "no" ## Use GPU or not
## (prefer "no" in the current version, because the GPU part may have bugs when integrated in this Python interface)
CPU_Part = 1.0
GPU_Part = 0.0
#################################################
#################################################
## Setting the physical system and numerical method
Symmetry = "equatorial-symmetry" ## Symmetry of System: choose equatorial-symmetry、no-symmetry、octant-symmetry
Equation_Class = "BSSN" ## Evolution Equation: choose "BSSN", "BSSN-EScalar", "BSSN-EM", "Z4C"
## If "BSSN-EScalar" is chosen, it is necessary to set other parameters below
Initial_Data_Method = "Ansorg-TwoPuncture" ## initial data method: choose "Ansorg-TwoPuncture", "Lousto-Analytical", "Cao-Analytical", "KerrSchild-Analytical"
Time_Evolution_Method = "runge-kutta-45" ## time evolution method: choose "runge-kutta-45"
Finite_Diffenence_Method = "4th-order" ## finite-difference method: choose "2nd-order", "4th-order", "6th-order", "8th-order"
Debug_NaN_Check = 0 ## enable NaN checks in compute_rhs_bssn: 0 (off) or 1 (on)
#################################################
#################################################
## Setting the time evolutionary information
Start_Evolution_Time = 0.0 ## start evolution time t0
Final_Evolution_Time = 100.0 ## final evolution time t1 (Reduced for quick test)
Check_Time = 10.0
Dump_Time = 10.0 ## time inteval dT for dumping binary data
D2_Dump_Time = 10.0 ## dump the ascii data for 2d surface after dT'
Analysis_Time = 1.0 ## dump the puncture position and GW psi4 after dT"
Evolution_Step_Number = 10000000 ## stop the calculation after the maximal step number
Courant_Factor = 0.5 ## Courant Factor
Dissipation = 0.15 ## Kreiss-Oliger Dissipation Strength
#################################################
#################################################
## Setting the grid structure
basic_grid_set = "Patch" ## grid structure: choose "Patch" or "Shell-Patch"
grid_center_set = "Cell" ## grid center: chose "Cell" or "Vertex"
grid_level = 7 ## total number of AMR grid levels (Reduced from 9)
static_grid_level = 4 ## number of AMR static grid levels (Reduced from 5)
moving_grid_level = grid_level - static_grid_level ## number of AMR moving grid levels
analysis_level = 0
refinement_level = 3 ## time refinement start from this grid level
largest_box_xyz_max = [320.0, 320.0, 320.0] ## scale of the largest box
## not ne cess ary to be cubic for "Patch" grid s tructure
## need to be a cubic box for "Shell-Patch" grid structure
largest_box_xyz_min = - numpy.array(largest_box_xyz_max)
static_grid_number = 48 ## grid points of each static AMR grid (in x direction) (Reduced from 96)
## (grid points in y and z directions are automatically adjusted)
moving_grid_number = 24 ## grid points of each moving AMR grid (Reduced from 48)
shell_grid_number = [32, 32, 100] ## grid points of Shell-Patch grid
## in (phi, theta, r) direction
devide_factor = 2.0 ## resolution between different grid levels dh0/dh1, only support 2.0 now
static_grid_type = 'Linear' ## AMR static grid structure , only supports "Linear"
moving_grid_type = 'Linear' ## AMR moving grid structure , only supports "Linear"
quarter_sphere_number = 48 ## grid number of 1/4 s pher ical surface (Reduced from 96)
## (which is needed for evaluating the spherical surface integral)
#################################################
#################################################
## Setting the puncture information
puncture_number = 2
position_BH = numpy.zeros( (puncture_number, 3) )
parameter_BH = numpy.zeros( (puncture_number, 3) )
dimensionless_spin_BH = numpy.zeros( (puncture_number, 3) )
momentum_BH = numpy.zeros( (puncture_number, 3) )
puncture_data_set = "Manually" ## Method to give Punctures positions and momentum
## choose "Manually" or "Automatically-BBH"
## Prefer to choose "Manually", because "Automatically-BBH" is developing now
## initial orbital distance and ellipticity for BBHs system
## ( needed for "Automatically-BBH" case , not affect the "Manually" case )
Distance = 10.0
e0 = 0.0
## black hole parameter (M Q* a*)
parameter_BH[0] = [ 36.0/(36.0+29.0), 0.0, +0.31 ]
parameter_BH[1] = [ 29.0/(36.0+29.0), 0.0, -0.46 ]
## dimensionless spin in each direction
dimensionless_spin_BH[0] = [ 0.0, 0.0, +0.31 ]
dimensionless_spin_BH[1] = [ 0.0, 0.0, -0.46 ]
## use Brugmann's convention
## -----0-----> y
## - +
#---------------------------------------------
## If puncture_data_set is chosen to be "Manually", it is necessary to set the position and momentum of each puncture manually
## initial position for each puncture
position_BH[0] = [ 0.0, 10.0*29.0/(36.0+29.0), 0.0 ]
position_BH[1] = [ 0.0, -10.0*36.0/(36.0+29.0), 0.0 ]
## initial mumentum for each puncture
## (needed for "Manually" case, does not affect the "Automatically-BBH" case)
momentum_BH[0] = [ -0.09530152296974252, -0.00084541526517121, 0.0 ]
momentum_BH[1] = [ +0.09530152296974252, +0.00084541526517121, 0.0 ]
#################################################
#################################################
## Setting the gravitational wave information
GW_L_max = 4 ## maximal L number in gravitational wave
GW_M_max = 4 ## maximal M number in gravitational wave
Detector_Number = 12 ## number of dector
Detector_Rmin = 50.0 ## nearest dector distance
Detector_Rmax = 160.0 ## farest dector distance
#################################################
#################################################
## Setting the apprent horizon
AHF_Find = "no" ## whether to find the apparent horizon: choose "yes" or "no"
AHF_Find_Every = 24
AHF_Dump_Time = 20.0
#################################################
#################################################
## Other parameters (testing)
## Only influence the Equation_Class = "BSSN-EScalar" case
FR_a2 = 3.0 ## f(R) = R + a2 * R^2
FR_l2 = 10000.0
FR_phi0 = 0.00005
FR_r0 = 120.0
FR_sigma0 = 8.0
FR_Choice = 2 ## Choice options: 1 2 3 4 5
## 1: phi(r) = phi0 * Exp(-(r-r0)**2/sigma0)
## V(r) = 0
## 2: phi(r) = phi0 * a2^2/(1+a2^2)
## V(r) = Exp(-8*Sqrt(PI/3)*phi(r)) * (1-Exp(4*Sqrt(PI/3)*phi(r)))**2 / (32*PI*a2)
## 3: Schrodinger-Newton gived by system phi(r)
## V(r) = Exp(-8*Sqrt(PI/3)*phi(r)) * (1-Exp(4*Sqrt(PI/3)*phi(r)))**2 / (32*PI*a2)
## 4: phi(r) = phi0 * 0.5 * ( tanh((r+r0)/sigma0) - tanh((r-r0)/sigma0) )
## V(r) = 0
## f(R) = R + a2*R^2 with a2 = +oo
## 5: phi(r) = phi0 * Exp(-(r-r0)**2/sigma)
## V(r) = 0
#################################################
#################################################
## Other parameters (testing)
## (please do not change if not necessary)
boundary_choice = "BAM-choice" ## Sommerfeld boundary condition : choose "BAM-choice" or "Shibata-choice"
## prefer "BAM-choice"
gauge_choice = 0 ## gauge choice
## 0: B^i gauge
## 1: David's puncture gauge
## 2: MB B^i gauge
## 3: RIT B^i gauge
## 4: MB beta gauge
## 5: RIT beta gauge
## 6: MGB1 B^i gauge
## 7: MGB2 B^i gauge
## prefer 0 or 1
tetrad_type = 2 ## tetradtype
## v:r; u: phi; w: theta
## v^a = (x,y,z)
## 0: orthonormal order: v,u,w
## v^a = (x,y,z)
## m = (phi - i theta)/sqrt(2)
## following Frans, Eq.(8) of PRD 75, 124018(2007)
## 1: orthonormal order: w,u,v
## m = (theta + i phi)/sqrt(2)
## following Sperhake, Eq.(3.2) of PRD 85, 124062(2012)
## 2: orthonormal order: v,u,w
## v_a = (x,y,z)
## m = (phi - i theta)/sqrt(2)
## following Frans, Eq.(8) of PRD 75, 124018(2007)
## this version recommend set to 2
## prefer 2
#################################################

224
AMSS_NCKU_MiniProgram.py Normal file
View File

@@ -0,0 +1,224 @@
##################################################################
##
## AMSS-NCKU Numerical Relativity Mini Test Program
## Author: Assistant (based on Xiaoqu's code)
## 2026/01/20
##
## This script runs a scaled-down version of the GW150914 test case
## suitable for laptop testing.
##
##################################################################
import os
import shutil
import sys
import time
# --- Context Manager for Input File Swapping ---
class InputFileSwapper:
def __init__(self, mini_file="AMSS_NCKU_Input_Mini.py", target_file="AMSS_NCKU_Input.py"):
self.mini_file = mini_file
self.target_file = target_file
self.backup_file = target_file + ".bak"
self.swapped = False
def __enter__(self):
print(f"[MiniProgram] Swapping {self.target_file} with {self.mini_file}...")
if os.path.exists(self.target_file):
shutil.move(self.target_file, self.backup_file)
shutil.copy(self.mini_file, self.target_file)
self.swapped = True
return self
def __exit__(self, exc_type, exc_value, traceback):
if self.swapped:
print(f"[MiniProgram] Restoring original {self.target_file}...")
os.remove(self.target_file)
if os.path.exists(self.backup_file):
shutil.move(self.backup_file, self.target_file)
def main():
# Use the swapper to ensure all imported modules see the mini configuration
with InputFileSwapper():
# Import modules AFTER swapping input file
try:
import AMSS_NCKU_Input as input_data
import print_information
import setup
import numerical_grid
import generate_macrodef
import makefile_and_run
import generate_TwoPuncture_input
import renew_puncture_parameter
import plot_xiaoqu
import plot_GW_strain_amplitude_xiaoqu
except ImportError as e:
print(f"Error importing modules: {e}")
return
print_information.print_program_introduction()
print("\n" + "#"*60)
print(" RUNNING MINI TEST CASE: GW150914-mini")
print("#"*60 + "\n")
# --- Directory Setup ---
File_directory = os.path.join(input_data.File_directory)
if os.path.exists(File_directory):
print(f" Output directory '{File_directory}' exists. Removing for mini test...")
shutil.rmtree(File_directory, ignore_errors=True)
os.mkdir(File_directory)
shutil.copy("AMSS_NCKU_Input.py", File_directory) # Copies the current (mini) input
output_directory = os.path.join(File_directory, "AMSS_NCKU_output")
os.mkdir(output_directory)
binary_results_directory = os.path.join(output_directory, input_data.Output_directory)
os.mkdir(binary_results_directory)
figure_directory = os.path.join(File_directory, "figure")
os.mkdir(figure_directory)
print(" Output directories generated.\n")
# --- Setup and Input Generation ---
setup.print_input_data(File_directory)
setup.generate_AMSSNCKU_input()
setup.print_puncture_information()
print("\n Generating AMSS-NCKU input parfile...")
numerical_grid.append_AMSSNCKU_cgh_input()
print("\n Plotting initial grid...")
numerical_grid.plot_initial_grid()
print("\n Generating macro files...")
generate_macrodef.generate_macrodef_h()
generate_macrodef.generate_macrodef_fh()
# --- Compilation Preparation ---
print("\n Preparing to compile and run...")
AMSS_NCKU_source_path = "AMSS_NCKU_source"
AMSS_NCKU_source_copy = os.path.join(File_directory, "AMSS_NCKU_source_copy")
if not os.path.exists(AMSS_NCKU_source_path):
print(" Error: AMSS_NCKU_source not found! Please run in the project root.")
return
shutil.copytree(AMSS_NCKU_source_path, AMSS_NCKU_source_copy)
macrodef_h_path = os.path.join(File_directory, "macrodef.h")
macrodef_fh_path = os.path.join(File_directory, "macrodef.fh")
shutil.copy2(macrodef_h_path, AMSS_NCKU_source_copy)
shutil.copy2(macrodef_fh_path, AMSS_NCKU_source_copy)
# --- Compilation ---
cwd = os.getcwd()
os.chdir(AMSS_NCKU_source_copy)
print(" Compiling ABE...")
makefile_and_run.makefile_ABE()
if (input_data.Initial_Data_Method == "Ansorg-TwoPuncture" ):
print(" Compiling TwoPunctureABE...")
makefile_and_run.makefile_TwoPunctureABE()
os.chdir(cwd)
# --- Copy Executables ---
if (input_data.GPU_Calculation == "no"):
ABE_file = os.path.join(AMSS_NCKU_source_copy, "ABE")
else:
ABE_file = os.path.join(AMSS_NCKU_source_copy, "ABEGPU")
if not os.path.exists(ABE_file):
print(" Error: ABE executable compilation failed.")
return
shutil.copy2(ABE_file, output_directory)
TwoPuncture_file = os.path.join(AMSS_NCKU_source_copy, "TwoPunctureABE")
if (input_data.Initial_Data_Method == "Ansorg-TwoPuncture" ):
if not os.path.exists(TwoPuncture_file):
print(" Error: TwoPunctureABE compilation failed.")
return
shutil.copy2(TwoPuncture_file, output_directory)
# --- Execution ---
start_time = time.time()
if (input_data.Initial_Data_Method == "Ansorg-TwoPuncture" ):
print("\n Generating TwoPuncture input...")
generate_TwoPuncture_input.generate_AMSSNCKU_TwoPuncture_input()
AMSS_NCKU_TwoPuncture_inputfile = 'AMSS-NCKU-TwoPuncture.input'
AMSS_NCKU_TwoPuncture_inputfile_path = os.path.join( File_directory, AMSS_NCKU_TwoPuncture_inputfile )
shutil.copy2( AMSS_NCKU_TwoPuncture_inputfile_path, os.path.join(output_directory, 'TwoPunctureinput.par') )
print(" Running TwoPunctureABE...")
os.chdir(output_directory)
makefile_and_run.run_TwoPunctureABE()
os.chdir(cwd)
# Update Puncture Parameter
renew_puncture_parameter.append_AMSSNCKU_BSSN_input(File_directory, output_directory)
AMSS_NCKU_inputfile = 'AMSS-NCKU.input'
AMSS_NCKU_inputfile_path = os.path.join(File_directory, AMSS_NCKU_inputfile)
shutil.copy2( AMSS_NCKU_inputfile_path, os.path.join(output_directory, 'input.par') )
print("\n Input files ready. Launching ABE...")
os.chdir(output_directory)
makefile_and_run.run_ABE()
os.chdir(cwd)
end_time = time.time()
elapsed_time = end_time - start_time
# --- Post-processing ---
print("\n Copying output files for inspection...")
AMSS_NCKU_error_file_path = os.path.join(binary_results_directory, "setting.par")
if os.path.exists(AMSS_NCKU_error_file_path):
shutil.copy( AMSS_NCKU_error_file_path, os.path.join(output_directory, "AMSSNCKU_setting_parameter") )
AMSS_NCKU_error_file_path = os.path.join(binary_results_directory, "Error.log")
if os.path.exists(AMSS_NCKU_error_file_path):
shutil.copy( AMSS_NCKU_error_file_path, os.path.join(output_directory, "Error.log") )
for fname in ["bssn_BH.dat", "bssn_ADMQs.dat", "bssn_psi4.dat", "bssn_constraint.dat"]:
fpath = os.path.join(binary_results_directory, fname)
if os.path.exists(fpath):
shutil.copy(fpath, os.path.join(output_directory, fname))
# --- Plotting ---
print("\n Plotting results...")
try:
plot_xiaoqu.generate_puncture_orbit_plot( binary_results_directory, figure_directory )
plot_xiaoqu.generate_puncture_orbit_plot3D( binary_results_directory, figure_directory )
plot_xiaoqu.generate_puncture_distence_plot( binary_results_directory, figure_directory )
for i in range(input_data.Detector_Number):
plot_xiaoqu.generate_gravitational_wave_psi4_plot( binary_results_directory, figure_directory, i )
plot_GW_strain_amplitude_xiaoqu.generate_gravitational_wave_amplitude_plot( binary_results_directory, figure_directory, i )
for i in range(input_data.Detector_Number):
plot_xiaoqu.generate_ADMmass_plot( binary_results_directory, figure_directory, i )
for i in range(input_data.grid_level):
plot_xiaoqu.generate_constraint_check_plot( binary_results_directory, figure_directory, i )
plot_xiaoqu.generate_binary_data_plot( binary_results_directory, figure_directory )
except Exception as e:
print(f"Warning: Plotting failed: {e}")
print(f"\n Program Cost = {elapsed_time:.2f} Seconds \n")
print(" AMSS-NCKU-Python simulation finished (Mini Test).\n")
if __name__ == "__main__":
main()

279
AMSS_NCKU_Verify_ASC26.py Normal file
View File

@@ -0,0 +1,279 @@
#!/usr/bin/env python3
"""
AMSS-NCKU GW150914 Simulation Regression Test Script
Verification Requirements:
1. XY-plane trajectory RMS error < 1% (Optimized vs. baseline, max of BH1 and BH2)
2. ADM constraint violation < 2 (Grid Level 0)
RMS Calculation Method:
- Computes trajectory deviation on the XY plane independently for BH1 and BH2
- For each black hole: RMS = sqrt((1/M) * sum((Δr_i / r_i^max)^2)) × 100%
- Final RMS = max(RMS_BH1, RMS_BH2)
Usage: python3 AMSS_NCKU_Verify_ASC26.py [output_dir]
Default: output_dir = GW150914/AMSS_NCKU_output
Reference: GW150914-origin (baseline simulation)
"""
import numpy as np
import sys
import os
# ANSI Color Codes
class Color:
GREEN = '\033[92m'
RED = '\033[91m'
YELLOW = '\033[93m'
BLUE = '\033[94m'
BOLD = '\033[1m'
RESET = '\033[0m'
def get_status_text(passed):
if passed:
return f"{Color.GREEN}{Color.BOLD}PASS{Color.RESET}"
else:
return f"{Color.RED}{Color.BOLD}FAIL{Color.RESET}"
def load_bh_trajectory(filepath):
"""Load black hole trajectory data"""
data = np.loadtxt(filepath)
return {
'time': data[:, 0],
'x1': data[:, 1], 'y1': data[:, 2], 'z1': data[:, 3],
'x2': data[:, 4], 'y2': data[:, 5], 'z2': data[:, 6]
}
def load_constraint_data(filepath):
"""Load constraint violation data"""
data = []
with open(filepath, 'r') as f:
for line in f:
if line.startswith('#'):
continue
parts = line.split()
if len(parts) >= 8:
data.append([float(x) for x in parts[:8]])
return np.array(data)
def calculate_rms_error(bh_data_ref, bh_data_target):
"""
Calculate trajectory-based RMS error on the XY plane between baseline and optimized simulations.
This function computes the RMS error independently for BH1 and BH2 trajectories,
then returns the maximum of the two as the final RMS error metric.
For each black hole, the RMS is calculated as:
RMS = sqrt( (1/M) * sum( (Δr_i / r_i^max)^2 ) ) × 100%
where:
Δr_i = sqrt((x_ref,i - x_new,i)^2 + (y_ref,i - y_new,i)^2)
r_i^max = max(sqrt(x_ref,i^2 + y_ref,i^2), sqrt(x_new,i^2 + y_new,i^2))
Args:
bh_data_ref: Reference (baseline) trajectory data
bh_data_target: Target (optimized) trajectory data
Returns:
rms_value: Final RMS error as a percentage (max of BH1 and BH2)
error: Error message if any
"""
# Align data: truncate to the length of the shorter dataset
M = min(len(bh_data_ref['time']), len(bh_data_target['time']))
if M < 10:
return None, "Insufficient data points for comparison"
# Extract XY coordinates for both black holes
x1_ref = bh_data_ref['x1'][:M]
y1_ref = bh_data_ref['y1'][:M]
x2_ref = bh_data_ref['x2'][:M]
y2_ref = bh_data_ref['y2'][:M]
x1_new = bh_data_target['x1'][:M]
y1_new = bh_data_target['y1'][:M]
x2_new = bh_data_target['x2'][:M]
y2_new = bh_data_target['y2'][:M]
# Calculate RMS for BH1
delta_r1 = np.sqrt((x1_ref - x1_new)**2 + (y1_ref - y1_new)**2)
r1_ref = np.sqrt(x1_ref**2 + y1_ref**2)
r1_new = np.sqrt(x1_new**2 + y1_new**2)
r1_max = np.maximum(r1_ref, r1_new)
# Calculate RMS for BH2
delta_r2 = np.sqrt((x2_ref - x2_new)**2 + (y2_ref - y2_new)**2)
r2_ref = np.sqrt(x2_ref**2 + y2_ref**2)
r2_new = np.sqrt(x2_new**2 + y2_new**2)
r2_max = np.maximum(r2_ref, r2_new)
# Avoid division by zero for BH1
valid_mask1 = r1_max > 1e-15
if np.sum(valid_mask1) < 10:
return None, "Insufficient valid data points for BH1"
terms1 = (delta_r1[valid_mask1] / r1_max[valid_mask1])**2
rms_bh1 = np.sqrt(np.mean(terms1)) * 100
# Avoid division by zero for BH2
valid_mask2 = r2_max > 1e-15
if np.sum(valid_mask2) < 10:
return None, "Insufficient valid data points for BH2"
terms2 = (delta_r2[valid_mask2] / r2_max[valid_mask2])**2
rms_bh2 = np.sqrt(np.mean(terms2)) * 100
# Final RMS is the maximum of BH1 and BH2
rms_final = max(rms_bh1, rms_bh2)
return rms_final, None
def analyze_constraint_violation(constraint_data, n_levels=9):
"""
Analyze ADM constraint violation
Return maximum constraint violation for Grid Level 0
"""
# Extract Grid Level 0 data (first entry for each time step)
level0_data = constraint_data[::n_levels]
# Calculate maximum absolute value for each constraint
results = {
'Ham': np.max(np.abs(level0_data[:, 1])),
'Px': np.max(np.abs(level0_data[:, 2])),
'Py': np.max(np.abs(level0_data[:, 3])),
'Pz': np.max(np.abs(level0_data[:, 4])),
'Gx': np.max(np.abs(level0_data[:, 5])),
'Gy': np.max(np.abs(level0_data[:, 6])),
'Gz': np.max(np.abs(level0_data[:, 7]))
}
results['max_violation'] = max(results.values())
return results
def print_header():
"""Print report header"""
print("\n" + Color.BLUE + Color.BOLD + "=" * 65 + Color.RESET)
print(Color.BOLD + " AMSS-NCKU GW150914 Simulation Regression Test Report" + Color.RESET)
print(Color.BLUE + Color.BOLD + "=" * 65 + Color.RESET)
def print_rms_results(rms_rel, error, threshold=1.0):
"""Print RMS error results"""
print(f"\n{Color.BOLD}1. RMS Error Analysis (Baseline vs Optimized){Color.RESET}")
print("-" * 45)
if error:
print(f" {Color.RED}Error: {error}{Color.RESET}")
return False
passed = rms_rel < threshold
print(f" RMS relative error: {rms_rel:.4f}%")
print(f" Requirement: < {threshold}%")
print(f" Status: {get_status_text(passed)}")
return passed
def print_constraint_results(results, threshold=2.0):
"""Print constraint violation results"""
print(f"\n{Color.BOLD}2. ADM Constraint Violation Analysis (Grid Level 0){Color.RESET}")
print("-" * 45)
names = ['Ham', 'Px', 'Py', 'Pz', 'Gx', 'Gy', 'Gz']
for i, name in enumerate(names):
print(f" Max |{name:3}|: {results[name]:.6f}", end=" ")
if (i + 1) % 2 == 0: print()
if len(names) % 2 != 0: print()
passed = results['max_violation'] < threshold
print(f"\n Maximum violation: {results['max_violation']:.6f}")
print(f" Requirement: < {threshold}")
print(f" Status: {get_status_text(passed)}")
return passed
def print_summary(rms_passed, constraint_passed):
"""Print summary"""
print("\n" + Color.BLUE + Color.BOLD + "=" * 65 + Color.RESET)
print(Color.BOLD + "Verification Summary" + Color.RESET)
print(Color.BLUE + Color.BOLD + "=" * 65 + Color.RESET)
all_passed = rms_passed and constraint_passed
res_rms = get_status_text(rms_passed)
res_con = get_status_text(constraint_passed)
print(f" [1] RMS trajectory check: {res_rms}")
print(f" [2] ADM constraint check: {res_con}")
final_status = f"{Color.GREEN}{Color.BOLD}ALL CHECKS PASSED{Color.RESET}" if all_passed else f"{Color.RED}{Color.BOLD}SOME CHECKS FAILED{Color.RESET}"
print(f"\n Overall result: {final_status}")
print(Color.BLUE + Color.BOLD + "=" * 65 + Color.RESET + "\n")
return all_passed
def main():
# Determine target (optimized) output directory
if len(sys.argv) > 1:
target_dir = sys.argv[1]
else:
script_dir = os.path.dirname(os.path.abspath(__file__))
target_dir = os.path.join(script_dir, "GW150914/AMSS_NCKU_output")
# Determine reference (baseline) directory
script_dir = os.path.dirname(os.path.abspath(__file__))
reference_dir = os.path.join(script_dir, "GW150914-origin/AMSS_NCKU_output")
# Data file paths
bh_file_ref = os.path.join(reference_dir, "bssn_BH.dat")
bh_file_target = os.path.join(target_dir, "bssn_BH.dat")
constraint_file = os.path.join(target_dir, "bssn_constraint.dat")
# Check if files exist
if not os.path.exists(bh_file_ref):
print(f"{Color.RED}{Color.BOLD}Error:{Color.RESET} Baseline trajectory file not found: {bh_file_ref}")
sys.exit(1)
if not os.path.exists(bh_file_target):
print(f"{Color.RED}{Color.BOLD}Error:{Color.RESET} Target trajectory file not found: {bh_file_target}")
sys.exit(1)
if not os.path.exists(constraint_file):
print(f"{Color.RED}{Color.BOLD}Error:{Color.RESET} Constraint data file not found: {constraint_file}")
sys.exit(1)
# Print header
print_header()
print(f"\n{Color.BOLD}Reference (Baseline):{Color.RESET} {Color.BLUE}{reference_dir}{Color.RESET}")
print(f"{Color.BOLD}Target (Optimized): {Color.RESET} {Color.BLUE}{target_dir}{Color.RESET}")
# Load data
bh_data_ref = load_bh_trajectory(bh_file_ref)
bh_data_target = load_bh_trajectory(bh_file_target)
constraint_data = load_constraint_data(constraint_file)
# Calculate RMS error
rms_rel, error = calculate_rms_error(bh_data_ref, bh_data_target)
rms_passed = print_rms_results(rms_rel, error)
# Analyze constraint violation
constraint_results = analyze_constraint_violation(constraint_data)
constraint_passed = print_constraint_results(constraint_results)
# Print summary
all_passed = print_summary(rms_passed, constraint_passed)
# Return exit code
sys.exit(0 if all_passed else 1)
if __name__ == "__main__":
main()

View File

@@ -37,57 +37,51 @@ close(77)
end program checkFFT end program checkFFT
#endif #endif
!-------------
! Optimized FFT using Intel oneMKL DFTI
! Mathematical equivalence: Standard DFT definition
! Forward (isign=1): X[k] = sum_{n=0}^{N-1} x[n] * exp(-2*pi*i*k*n/N)
! Backward (isign=-1): X[k] = sum_{n=0}^{N-1} x[n] * exp(+2*pi*i*k*n/N)
! Input/Output: dataa is interleaved complex array [Re(0),Im(0),Re(1),Im(1),...]
!------------- !-------------
SUBROUTINE four1(dataa,nn,isign) SUBROUTINE four1(dataa,nn,isign)
use MKL_DFTI
implicit none implicit none
INTEGER::isign,nn INTEGER, intent(in) :: isign, nn
double precision,dimension(2*nn)::dataa DOUBLE PRECISION, dimension(2*nn), intent(inout) :: dataa
INTEGER::i,istep,j,m,mmax,n
double precision::tempi,tempr type(DFTI_DESCRIPTOR), pointer :: desc
DOUBLE PRECISION::theta,wi,wpi,wpr,wr,wtemp integer :: status
n=2*nn
j=1 ! Create DFTI descriptor for 1D complex-to-complex transform
do i=1,n,2 status = DftiCreateDescriptor(desc, DFTI_DOUBLE, DFTI_COMPLEX, 1, nn)
if(j.gt.i)then if (status /= 0) return
tempr=dataa(j)
tempi=dataa(j+1) ! Set input/output storage as interleaved complex (default)
dataa(j)=dataa(i) status = DftiSetValue(desc, DFTI_PLACEMENT, DFTI_INPLACE)
dataa(j+1)=dataa(i+1) if (status /= 0) then
dataa(i)=tempr status = DftiFreeDescriptor(desc)
dataa(i+1)=tempi return
endif
m=nn
1 if ((m.ge.2).and.(j.gt.m)) then
j=j-m
m=m/2
goto 1
endif
j=j+m
enddo
mmax=2
2 if (n.gt.mmax) then
istep=2*mmax
theta=6.28318530717959d0/(isign*mmax)
wpr=-2.d0*sin(0.5d0*theta)**2
wpi=sin(theta)
wr=1.d0
wi=0.d0
do m=1,mmax,2
do i=m,n,istep
j=i+mmax
tempr=sngl(wr)*dataa(j)-sngl(wi)*dataa(j+1)
tempi=sngl(wr)*dataa(j+1)+sngl(wi)*dataa(j)
dataa(j)=dataa(i)-tempr
dataa(j+1)=dataa(i+1)-tempi
dataa(i)=dataa(i)+tempr
dataa(i+1)=dataa(i+1)+tempi
enddo
wtemp=wr
wr=wr*wpr-wi*wpi+wr
wi=wi*wpr+wtemp*wpi+wi
enddo
mmax=istep
goto 2
endif endif
! Commit the descriptor
status = DftiCommitDescriptor(desc)
if (status /= 0) then
status = DftiFreeDescriptor(desc)
return
endif
! Execute FFT based on direction
if (isign == 1) then
! Forward FFT: exp(-2*pi*i*k*n/N)
status = DftiComputeForward(desc, dataa)
else
! Backward FFT: exp(+2*pi*i*k*n/N)
status = DftiComputeBackward(desc, dataa)
endif
! Free descriptor
status = DftiFreeDescriptor(desc)
return return
END SUBROUTINE four1 END SUBROUTINE four1

View File

@@ -27,6 +27,7 @@ using namespace std;
#endif #endif
#include "TwoPunctures.h" #include "TwoPunctures.h"
#include <mkl_cblas.h>
TwoPunctures::TwoPunctures(double mp, double mm, double b, TwoPunctures::TwoPunctures(double mp, double mm, double b,
double P_plusx, double P_plusy, double P_plusz, double P_plusx, double P_plusy, double P_plusz,
@@ -891,25 +892,17 @@ double TwoPunctures::norm1(double *v, int n)
/* -------------------------------------------------------------------------*/ /* -------------------------------------------------------------------------*/
double TwoPunctures::norm2(double *v, int n) double TwoPunctures::norm2(double *v, int n)
{ {
int i; // Optimized with oneMKL BLAS DNRM2
double result = 0; // Computes: sqrt(sum(v[i]^2))
return cblas_dnrm2(n, v, 1);
for (i = 0; i < n; i++)
result += v[i] * v[i];
return sqrt(result);
} }
/* -------------------------------------------------------------------------*/ /* -------------------------------------------------------------------------*/
double TwoPunctures::scalarproduct(double *v, double *w, int n) double TwoPunctures::scalarproduct(double *v, double *w, int n)
{ {
int i; // Optimized with oneMKL BLAS DDOT
double result = 0; // Computes: sum(v[i] * w[i])
return cblas_ddot(n, v, 1, w, 1);
for (i = 0; i < n; i++)
result += v[i] * w[i];
return result;
} }
/* -------------------------------------------------------------------------*/ /* -------------------------------------------------------------------------*/

View File

@@ -18,7 +18,7 @@ using namespace std;
#include <fstream> #include <fstream>
#endif #endif
static void compare_result_gpu(int ftag1,double * datac,int data_num){ void compare_result_gpu(int ftag1,double * datac,int data_num){
double * data = (double*)malloc(sizeof(double)*data_num); double * data = (double*)malloc(sizeof(double)*data_num);
cudaMemcpy(data, datac, data_num * sizeof(double), cudaMemcpyDeviceToHost); cudaMemcpy(data, datac, data_num * sizeof(double), cudaMemcpyDeviceToHost);
compare_result(ftag1,data,data_num); compare_result(ftag1,data,data_num);
@@ -83,7 +83,7 @@ inline void sub_enforce_ga(int matrix_size){
double * trA = M_ chin1; double * trA = M_ chin1;
enforce_ga<<<GRID_DIM,BLOCK_DIM>>>(trA); enforce_ga<<<GRID_DIM,BLOCK_DIM>>>(trA);
cudaMemset(trA,0,matrix_size * sizeof(double)); cudaMemset(trA,0,matrix_size * sizeof(double));
cudaDeviceSynchronize(); cudaThreadSynchronize();
//cudaMemset(Mh_ gupxx,0,matrix_size * sizeof(double)); //cudaMemset(Mh_ gupxx,0,matrix_size * sizeof(double));
//trA gxx,gyy,gzz gupxx,gupxy,gupxz,gupyy,gupyz,gupzz //trA gxx,gyy,gzz gupxx,gupxy,gupxz,gupyy,gupyz,gupzz
@@ -273,13 +273,13 @@ __global__ void sub_symmetry_bd_partK(int ord,double * func, double * funcc,doub
#endif //ifdef Vertex #endif //ifdef Vertex
inline void sub_symmetry_bd(int ord,double * func, double * funcc,double * SoA){ inline void sub_symmetry_bd(int ord,double * func, double * funcc,double * SoA){
sub_symmetry_bd_partF<<<GRID_DIM,BLOCK_DIM>>>(ord,func,funcc); sub_symmetry_bd_partF<<<GRID_DIM,BLOCK_DIM>>>(ord,func,funcc);
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_symmetry_bd_partI<<<GRID_DIM,BLOCK_DIM>>>(ord,func,funcc,SoA[0]); sub_symmetry_bd_partI<<<GRID_DIM,BLOCK_DIM>>>(ord,func,funcc,SoA[0]);
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_symmetry_bd_partJ<<<GRID_DIM,BLOCK_DIM>>>(ord,func,funcc,SoA[1]); sub_symmetry_bd_partJ<<<GRID_DIM,BLOCK_DIM>>>(ord,func,funcc,SoA[1]);
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_symmetry_bd_partK<<<GRID_DIM,BLOCK_DIM>>>(ord,func,funcc,SoA[2]); sub_symmetry_bd_partK<<<GRID_DIM,BLOCK_DIM>>>(ord,func,funcc,SoA[2]);
cudaDeviceSynchronize(); cudaThreadSynchronize();
} }
@@ -378,9 +378,9 @@ inline void sub_fdderivs(double * f,double *fh,double *fxx,double *fxy,double *f
cudaMemset(fyy,0,_3D_SIZE[0] * sizeof(double)); cudaMemset(fyy,0,_3D_SIZE[0] * sizeof(double));
cudaMemset(fyz,0,_3D_SIZE[0] * sizeof(double)); cudaMemset(fyz,0,_3D_SIZE[0] * sizeof(double));
cudaMemset(fzz,0,_3D_SIZE[0] * sizeof(double)); cudaMemset(fzz,0,_3D_SIZE[0] * sizeof(double));
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_fdderivs_part1<<<GRID_DIM,BLOCK_DIM>>>(f,fh,fxx,fxy,fxz,fyy,fyz,fzz); sub_fdderivs_part1<<<GRID_DIM,BLOCK_DIM>>>(f,fh,fxx,fxy,fxz,fyy,fyz,fzz);
cudaDeviceSynchronize(); cudaThreadSynchronize();
} }
__global__ void sub_fderivs_part1(double * f,double * fh,double *fx,double *fy,double *fz ) __global__ void sub_fderivs_part1(double * f,double * fh,double *fx,double *fy,double *fz )
@@ -445,9 +445,9 @@ inline void sub_fderivs(double * f,double * fh,double *fx,double *fy,double *fz,
cudaMemset(fy,0,_3D_SIZE[0] * sizeof(double)); cudaMemset(fy,0,_3D_SIZE[0] * sizeof(double));
cudaMemset(fz,0,_3D_SIZE[0] * sizeof(double)); cudaMemset(fz,0,_3D_SIZE[0] * sizeof(double));
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_fderivs_part1<<<GRID_DIM,BLOCK_DIM>>>(f,fh,fx,fy,fz); sub_fderivs_part1<<<GRID_DIM,BLOCK_DIM>>>(f,fh,fx,fy,fz);
cudaDeviceSynchronize(); cudaThreadSynchronize();
} }
__global__ void computeRicci_part1(double * dst) __global__ void computeRicci_part1(double * dst)
@@ -465,9 +465,9 @@ __global__ void computeRicci_part1(double * dst)
inline void computeRicci(double * src,double* dst,double * SoA, Meta* meta) inline void computeRicci(double * src,double* dst,double * SoA, Meta* meta)
{ {
sub_fdderivs(src,Mh_ fh,Mh_ fxx,Mh_ fxy,Mh_ fxz,Mh_ fyy,Mh_ fyz,Mh_ fzz,SoA); sub_fdderivs(src,Mh_ fh,Mh_ fxx,Mh_ fxy,Mh_ fxz,Mh_ fyy,Mh_ fyz,Mh_ fzz,SoA);
cudaDeviceSynchronize(); cudaThreadSynchronize();
computeRicci_part1<<<GRID_DIM,BLOCK_DIM>>>(dst); computeRicci_part1<<<GRID_DIM,BLOCK_DIM>>>(dst);
cudaDeviceSynchronize(); cudaThreadSynchronize();
}/*Exception*/ }/*Exception*/
@@ -524,9 +524,9 @@ __global__ void sub_kodis_part1(double *f,double *fh,double *f_rhs)
inline void sub_kodis(double *f,double *fh,double *f_rhs,double *SoA) inline void sub_kodis(double *f,double *fh,double *f_rhs,double *SoA)
{ {
sub_symmetry_bd(3,f,fh,SoA); sub_symmetry_bd(3,f,fh,SoA);
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_kodis_part1<<<GRID_DIM,BLOCK_DIM>>>(f,fh,f_rhs); sub_kodis_part1<<<GRID_DIM,BLOCK_DIM>>>(f,fh,f_rhs);
cudaDeviceSynchronize(); cudaThreadSynchronize();
} }
__global__ void sub_lopsided_part1(double *f,double* fh,double *f_rhs,double *Sfx,double *Sfy,double *Sfz) __global__ void sub_lopsided_part1(double *f,double* fh,double *f_rhs,double *Sfx,double *Sfy,double *Sfz)
@@ -617,9 +617,9 @@ __global__ void sub_lopsided_part1(double *f,double* fh,double *f_rhs,double *S
inline void sub_lopsided(double *f,double*fh,double *f_rhs,double *Sfx,double *Sfy,double *Sfz,double *SoA){ inline void sub_lopsided(double *f,double*fh,double *f_rhs,double *Sfx,double *Sfy,double *Sfz,double *SoA){
sub_symmetry_bd(3,f,fh,SoA); sub_symmetry_bd(3,f,fh,SoA);
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_lopsided_part1<<<GRID_DIM,BLOCK_DIM>>>(f,fh,f_rhs,Sfx,Sfy,Sfz); sub_lopsided_part1<<<GRID_DIM,BLOCK_DIM>>>(f,fh,f_rhs,Sfx,Sfy,Sfz);
cudaDeviceSynchronize(); cudaThreadSynchronize();
} }
__global__ void compute_rhs_bssn_part1() __global__ void compute_rhs_bssn_part1()
@@ -2656,13 +2656,13 @@ int gpu_rhs(int calledby, int mpi_rank, int *ex, double &T,double *X, double *Y,
#ifdef TIMING1 #ifdef TIMING1
cudaDeviceSynchronize(); cudaThreadSynchronize();
gettimeofday(&tv2, NULL); gettimeofday(&tv2, NULL);
cout<<"TIME USED"<<TimeBetween(tv1, tv2)<<endl; cout<<"TIME USED"<<TimeBetween(tv1, tv2)<<endl;
#endif #endif
//cout<<"GPU meta data ready.\n"; //cout<<"GPU meta data ready.\n";
cudaDeviceSynchronize(); cudaThreadSynchronize();
//--------------test constant memory address & value-------------- //--------------test constant memory address & value--------------
/* double rank = mpi_rank; /* double rank = mpi_rank;
@@ -2685,7 +2685,7 @@ int gpu_rhs(int calledby, int mpi_rank, int *ex, double &T,double *X, double *Y,
//sub_enforce_ga(matrix_size); //sub_enforce_ga(matrix_size);
//4.1-----compute rhs--------- //4.1-----compute rhs---------
compute_rhs_bssn_part1<<<GRID_DIM,BLOCK_DIM>>>(); compute_rhs_bssn_part1<<<GRID_DIM,BLOCK_DIM>>>();
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_fderivs(Mh_ betax,Mh_ fh,Mh_ betaxx,Mh_ betaxy,Mh_ betaxz,ass); sub_fderivs(Mh_ betax,Mh_ fh,Mh_ betaxx,Mh_ betaxy,Mh_ betaxz,ass);
sub_fderivs(Mh_ betay,Mh_ fh,Mh_ betayx,Mh_ betayy,Mh_ betayz,sas); sub_fderivs(Mh_ betay,Mh_ fh,Mh_ betayx,Mh_ betayy,Mh_ betayz,sas);
@@ -2701,7 +2701,7 @@ int gpu_rhs(int calledby, int mpi_rank, int *ex, double &T,double *X, double *Y,
sub_fderivs(Mh_ gyz,Mh_ fh,Mh_ gyzx,Mh_ gyzy,Mh_ gyzz, saa); sub_fderivs(Mh_ gyz,Mh_ fh,Mh_ gyzx,Mh_ gyzy,Mh_ gyzz, saa);
compute_rhs_bssn_part2<<<GRID_DIM,BLOCK_DIM>>>(); compute_rhs_bssn_part2<<<GRID_DIM,BLOCK_DIM>>>();
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_fdderivs(Mh_ betax,Mh_ fh,Mh_ gxxx,Mh_ gxyx,Mh_ gxzx,Mh_ gyyx,Mh_ gyzx,Mh_ gzzx,ass); sub_fdderivs(Mh_ betax,Mh_ fh,Mh_ gxxx,Mh_ gxyx,Mh_ gxzx,Mh_ gyyx,Mh_ gyzx,Mh_ gzzx,ass);
sub_fdderivs(Mh_ betay,Mh_ fh,Mh_ gxxy,Mh_ gxyy,Mh_ gxzy,Mh_ gyyy,Mh_ gyzy,Mh_ gzzy,sas); sub_fdderivs(Mh_ betay,Mh_ fh,Mh_ gxxy,Mh_ gxyy,Mh_ gxzy,Mh_ gyyy,Mh_ gyzy,Mh_ gzzy,sas);
@@ -2711,7 +2711,7 @@ int gpu_rhs(int calledby, int mpi_rank, int *ex, double &T,double *X, double *Y,
sub_fderivs( Mh_ Gamz, Mh_ fh,Mh_ Gamzx, Mh_ Gamzy, Mh_ Gamzz,ssa); sub_fderivs( Mh_ Gamz, Mh_ fh,Mh_ Gamzx, Mh_ Gamzy, Mh_ Gamzz,ssa);
compute_rhs_bssn_part3<<<GRID_DIM,BLOCK_DIM>>>(); compute_rhs_bssn_part3<<<GRID_DIM,BLOCK_DIM>>>();
cudaDeviceSynchronize(); cudaThreadSynchronize();
computeRicci(Mh_ dxx,Mh_ Rxx,sss, meta); computeRicci(Mh_ dxx,Mh_ Rxx,sss, meta);
computeRicci(Mh_ dyy,Mh_ Ryy,sss, meta); computeRicci(Mh_ dyy,Mh_ Ryy,sss, meta);
@@ -2720,20 +2720,20 @@ int gpu_rhs(int calledby, int mpi_rank, int *ex, double &T,double *X, double *Y,
computeRicci(Mh_ gxz,Mh_ Rxz,asa, meta); computeRicci(Mh_ gxz,Mh_ Rxz,asa, meta);
computeRicci(Mh_ gyz,Mh_ Ryz,saa, meta); computeRicci(Mh_ gyz,Mh_ Ryz,saa, meta);
cudaDeviceSynchronize(); cudaThreadSynchronize();
compute_rhs_bssn_part4<<<GRID_DIM,BLOCK_DIM>>>(); compute_rhs_bssn_part4<<<GRID_DIM,BLOCK_DIM>>>();
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_fdderivs(Mh_ chi,Mh_ fh,Mh_ fxx,Mh_ fxy,Mh_ fxz,Mh_ fyy,Mh_ fyz,Mh_ fzz,sss); sub_fdderivs(Mh_ chi,Mh_ fh,Mh_ fxx,Mh_ fxy,Mh_ fxz,Mh_ fyy,Mh_ fyz,Mh_ fzz,sss);
compute_rhs_bssn_part5<<<GRID_DIM,BLOCK_DIM>>>(); compute_rhs_bssn_part5<<<GRID_DIM,BLOCK_DIM>>>();
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_fdderivs(Mh_ Lap,Mh_ fh,Mh_ fxx,Mh_ fxy,Mh_ fxz,Mh_ fyy,Mh_ fyz,Mh_ fzz,sss); sub_fdderivs(Mh_ Lap,Mh_ fh,Mh_ fxx,Mh_ fxy,Mh_ fxz,Mh_ fyy,Mh_ fyz,Mh_ fzz,sss);
compute_rhs_bssn_part6<<<GRID_DIM,BLOCK_DIM>>>(); compute_rhs_bssn_part6<<<GRID_DIM,BLOCK_DIM>>>();
cudaDeviceSynchronize(); cudaThreadSynchronize();
#if (GAUGE == 2 || GAUGE == 3 || GAUGE == 4 || GAUGE == 5) #if (GAUGE == 2 || GAUGE == 3 || GAUGE == 4 || GAUGE == 5)
sub_fderivs(Mh_ chi,Mh_ fh, Mh_ dtSfx_rhs, Mh_ dtSfy_rhs, Mh_ dtSfz_rhs,sss); sub_fderivs(Mh_ chi,Mh_ fh, Mh_ dtSfx_rhs, Mh_ dtSfy_rhs, Mh_ dtSfz_rhs,sss);
@@ -2805,7 +2805,7 @@ int gpu_rhs(int calledby, int mpi_rank, int *ex, double &T,double *X, double *Y,
if(co == 0){ if(co == 0){
compute_rhs_bssn_part7<<<GRID_DIM,BLOCK_DIM>>>(); compute_rhs_bssn_part7<<<GRID_DIM,BLOCK_DIM>>>();
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_fderivs(Mh_ Axx,Mh_ fh,Mh_ gxxx,Mh_ gxxy,Mh_ gxxz,sss); sub_fderivs(Mh_ Axx,Mh_ fh,Mh_ gxxx,Mh_ gxxy,Mh_ gxxz,sss);
sub_fderivs(Mh_ Axy,Mh_ fh,Mh_ gxyx,Mh_ gxyy,Mh_ gxyz,aas); sub_fderivs(Mh_ Axy,Mh_ fh,Mh_ gxyx,Mh_ gxyy,Mh_ gxyz,aas);
@@ -2814,7 +2814,7 @@ int gpu_rhs(int calledby, int mpi_rank, int *ex, double &T,double *X, double *Y,
sub_fderivs(Mh_ Ayz,Mh_ fh,Mh_ gyzx,Mh_ gyzy,Mh_ gyzz,saa); sub_fderivs(Mh_ Ayz,Mh_ fh,Mh_ gyzx,Mh_ gyzy,Mh_ gyzz,saa);
sub_fderivs(Mh_ Azz,Mh_ fh,Mh_ gzzx,Mh_ gzzy,Mh_ gzzz,sss); sub_fderivs(Mh_ Azz,Mh_ fh,Mh_ gzzx,Mh_ gzzy,Mh_ gzzz,sss);
compute_rhs_bssn_part8<<<GRID_DIM,BLOCK_DIM>>>(); compute_rhs_bssn_part8<<<GRID_DIM,BLOCK_DIM>>>();
cudaDeviceSynchronize(); cudaThreadSynchronize();
} }
#if (ABV == 1) #if (ABV == 1)
@@ -2895,7 +2895,7 @@ int gpu_rhs(int calledby, int mpi_rank, int *ex, double &T,double *X, double *Y,
//-------------------FOR GPU TEST---------------------- //-------------------FOR GPU TEST----------------------
//----------------------------------------------------- //-----------------------------------------------------
#ifdef TIMING #ifdef TIMING
cudaDeviceSynchronize(); cudaThreadSynchronize();
gettimeofday(&tv2, NULL); gettimeofday(&tv2, NULL);
cout<<"MPI rank is: "<<mpi_rank<<" GPU TIME is"<<TimeBetween(tv1, tv2)<<" (s)."<<endl; cout<<"MPI rank is: "<<mpi_rank<<" GPU TIME is"<<TimeBetween(tv1, tv2)<<" (s)."<<endl;
#endif #endif

View File

@@ -4,17 +4,6 @@
#include "bssn_macro.h" #include "bssn_macro.h"
#include "macrodef.fh" #include "macrodef.fh"
// CUDA error checking macro for CUDA 13 compatibility
#define CUDA_SAFE_CALL(call) \
do { \
cudaError_t err = call; \
if (err != cudaSuccess) { \
fprintf(stderr, "CUDA error in %s:%d: %s\n", __FILE__, __LINE__, \
cudaGetErrorString(err)); \
exit(EXIT_FAILURE); \
} \
} while(0)
#define DEVICE_ID 0 #define DEVICE_ID 0
// #define DEVICE_ID_BY_MPI_RANK // #define DEVICE_ID_BY_MPI_RANK
#define GRID_DIM 256 #define GRID_DIM 256

View File

@@ -2134,9 +2134,7 @@ void bssn_class::Evolve(int Steps)
CheckPoint->write_Black_Hole_position(BH_num_input, BH_num, Porg0, Porgbr, Mass); CheckPoint->write_Black_Hole_position(BH_num_input, BH_num, Porg0, Porgbr, Mass);
CheckPoint->writecheck_cgh(PhysTime, GH); CheckPoint->writecheck_cgh(PhysTime, GH);
#ifdef WithShell
CheckPoint->writecheck_sh(PhysTime, SH); CheckPoint->writecheck_sh(PhysTime, SH);
#endif
CheckPoint->write_bssn(LastDump, Last2dDump, LastAnas); CheckPoint->write_bssn(LastDump, Last2dDump, LastAnas);
} }
} }

View File

@@ -20,7 +20,7 @@ using namespace std;
__device__ volatile unsigned int global_count = 0; __device__ volatile unsigned int global_count = 0;
static void compare_result_gpu(int ftag1,double * datac,int data_num){ void compare_result_gpu(int ftag1,double * datac,int data_num){
double * data = (double*)malloc(sizeof(double)*data_num); double * data = (double*)malloc(sizeof(double)*data_num);
cudaMemcpy(data, datac, data_num * sizeof(double), cudaMemcpyDeviceToHost); cudaMemcpy(data, datac, data_num * sizeof(double), cudaMemcpyDeviceToHost);
compare_result(ftag1,data,data_num); compare_result(ftag1,data,data_num);
@@ -153,11 +153,11 @@ __global__ void sub_symmetry_bd_ss_partJ(int ord,double * func, double * funcc,d
inline void sub_symmetry_bd_ss(int ord,double * func, double * funcc,double * SoA){ inline void sub_symmetry_bd_ss(int ord,double * func, double * funcc,double * SoA){
sub_symmetry_bd_ss_partF<<<GRID_DIM,BLOCK_DIM>>>(ord,func,funcc); sub_symmetry_bd_ss_partF<<<GRID_DIM,BLOCK_DIM>>>(ord,func,funcc);
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_symmetry_bd_ss_partI<<<GRID_DIM,BLOCK_DIM>>>(ord,func,funcc,SoA[0]); sub_symmetry_bd_ss_partI<<<GRID_DIM,BLOCK_DIM>>>(ord,func,funcc,SoA[0]);
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_symmetry_bd_ss_partJ<<<GRID_DIM,BLOCK_DIM>>>(ord,func,funcc,SoA[1]); sub_symmetry_bd_ss_partJ<<<GRID_DIM,BLOCK_DIM>>>(ord,func,funcc,SoA[1]);
cudaDeviceSynchronize(); cudaThreadSynchronize();
} }
__global__ void sub_fderivs_shc_part1(double *fx,double *fy,double *fz){ __global__ void sub_fderivs_shc_part1(double *fx,double *fy,double *fz){
@@ -247,13 +247,13 @@ inline void sub_fderivs_shc(int& sst,double * f,double * fh,double *fx,double *f
//cudaMemset(Msh_ gy,0,h_3D_SIZE[0] * sizeof(double)); //cudaMemset(Msh_ gy,0,h_3D_SIZE[0] * sizeof(double));
//cudaMemset(Msh_ gz,0,h_3D_SIZE[0] * sizeof(double)); //cudaMemset(Msh_ gz,0,h_3D_SIZE[0] * sizeof(double));
sub_symmetry_bd_ss(2,f,fh,SoA1); sub_symmetry_bd_ss(2,f,fh,SoA1);
cudaDeviceSynchronize(); cudaThreadSynchronize();
//compare_result_gpu(0,fh,h_3D_SIZE[2]); //compare_result_gpu(0,fh,h_3D_SIZE[2]);
sub_fderivs_sh<<<GRID_DIM,BLOCK_DIM>>>(fh,Msh_ gx,Msh_ gy,Msh_ gz); sub_fderivs_sh<<<GRID_DIM,BLOCK_DIM>>>(fh,Msh_ gx,Msh_ gy,Msh_ gz);
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_fderivs_shc_part1<<<GRID_DIM,BLOCK_DIM>>>(fx,fy,fz); sub_fderivs_shc_part1<<<GRID_DIM,BLOCK_DIM>>>(fx,fy,fz);
cudaDeviceSynchronize(); cudaThreadSynchronize();
//compare_result_gpu(1,fx,h_3D_SIZE[0]); //compare_result_gpu(1,fx,h_3D_SIZE[0]);
//compare_result_gpu(2,fy,h_3D_SIZE[0]); //compare_result_gpu(2,fy,h_3D_SIZE[0]);
//compare_result_gpu(3,fz,h_3D_SIZE[0]); //compare_result_gpu(3,fz,h_3D_SIZE[0]);
@@ -451,17 +451,17 @@ inline void sub_fdderivs_shc(int& sst,double * f,double * fh,
//fderivs_sh //fderivs_sh
sub_symmetry_bd_ss(2,f,fh,SoA1); sub_symmetry_bd_ss(2,f,fh,SoA1);
cudaDeviceSynchronize(); cudaThreadSynchronize();
//compare_result_gpu(1,fh,h_3D_SIZE[2]); //compare_result_gpu(1,fh,h_3D_SIZE[2]);
sub_fderivs_sh<<<GRID_DIM,BLOCK_DIM>>>(fh,Msh_ gx,Msh_ gy,Msh_ gz); sub_fderivs_sh<<<GRID_DIM,BLOCK_DIM>>>(fh,Msh_ gx,Msh_ gy,Msh_ gz);
cudaDeviceSynchronize(); cudaThreadSynchronize();
//fdderivs_sh //fdderivs_sh
sub_symmetry_bd_ss(2,f,fh,SoA1); sub_symmetry_bd_ss(2,f,fh,SoA1);
cudaDeviceSynchronize(); cudaThreadSynchronize();
//compare_result_gpu(21,fh,h_3D_SIZE[2]); //compare_result_gpu(21,fh,h_3D_SIZE[2]);
sub_fdderivs_sh<<<GRID_DIM,BLOCK_DIM>>>(fh,Msh_ gxx,Msh_ gxy,Msh_ gxz,Msh_ gyy,Msh_ gyz,Msh_ gzz); sub_fdderivs_sh<<<GRID_DIM,BLOCK_DIM>>>(fh,Msh_ gxx,Msh_ gxy,Msh_ gxz,Msh_ gyy,Msh_ gyz,Msh_ gzz);
cudaDeviceSynchronize(); cudaThreadSynchronize();
/*compare_result_gpu(11,Msh_ gx,h_3D_SIZE[0]); /*compare_result_gpu(11,Msh_ gx,h_3D_SIZE[0]);
compare_result_gpu(12,Msh_ gy,h_3D_SIZE[0]); compare_result_gpu(12,Msh_ gy,h_3D_SIZE[0]);
compare_result_gpu(13,Msh_ gz,h_3D_SIZE[0]); compare_result_gpu(13,Msh_ gz,h_3D_SIZE[0]);
@@ -472,7 +472,7 @@ inline void sub_fdderivs_shc(int& sst,double * f,double * fh,
compare_result_gpu(5,Msh_ gyz,h_3D_SIZE[0]); compare_result_gpu(5,Msh_ gyz,h_3D_SIZE[0]);
compare_result_gpu(6,Msh_ gzz,h_3D_SIZE[0]);*/ compare_result_gpu(6,Msh_ gzz,h_3D_SIZE[0]);*/
sub_fdderivs_shc_part1<<<GRID_DIM,BLOCK_DIM>>>(fxx,fxy,fxz,fyy,fyz,fzz); sub_fdderivs_shc_part1<<<GRID_DIM,BLOCK_DIM>>>(fxx,fxy,fxz,fyy,fyz,fzz);
cudaDeviceSynchronize(); cudaThreadSynchronize();
/*compare_result_gpu(1,fxx,h_3D_SIZE[0]); /*compare_result_gpu(1,fxx,h_3D_SIZE[0]);
compare_result_gpu(2,fxy,h_3D_SIZE[0]); compare_result_gpu(2,fxy,h_3D_SIZE[0]);
compare_result_gpu(3,fxz,h_3D_SIZE[0]); compare_result_gpu(3,fxz,h_3D_SIZE[0]);
@@ -496,9 +496,9 @@ __global__ void computeRicci_ss_part1(double * dst)
inline void computeRicci_ss(int &sst,double * src,double* dst,double * SoA, Meta* meta) inline void computeRicci_ss(int &sst,double * src,double* dst,double * SoA, Meta* meta)
{ {
sub_fdderivs_shc(sst,src,Mh_ fh,Mh_ fxx,Mh_ fxy,Mh_ fxz,Mh_ fyy,Mh_ fyz,Mh_ fzz,SoA); sub_fdderivs_shc(sst,src,Mh_ fh,Mh_ fxx,Mh_ fxy,Mh_ fxz,Mh_ fyy,Mh_ fyz,Mh_ fzz,SoA);
cudaDeviceSynchronize(); cudaThreadSynchronize();
computeRicci_ss_part1<<<GRID_DIM,BLOCK_DIM>>>(dst); computeRicci_ss_part1<<<GRID_DIM,BLOCK_DIM>>>(dst);
cudaDeviceSynchronize(); cudaThreadSynchronize();
} }
__global__ void sub_lopsided_ss_part1(double * dst) __global__ void sub_lopsided_ss_part1(double * dst)
@@ -516,9 +516,9 @@ __global__ void sub_lopsided_ss_part1(double * dst)
inline void sub_lopsided_ss(int& sst,double *src,double* dst,double *SoA) inline void sub_lopsided_ss(int& sst,double *src,double* dst,double *SoA)
{ {
sub_fderivs_shc(sst,src,Mh_ fh,Mh_ fxx,Mh_ fxy,Mh_ fxz,SoA); sub_fderivs_shc(sst,src,Mh_ fh,Mh_ fxx,Mh_ fxy,Mh_ fxz,SoA);
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_lopsided_ss_part1<<<GRID_DIM,BLOCK_DIM>>>(dst); sub_lopsided_ss_part1<<<GRID_DIM,BLOCK_DIM>>>(dst);
cudaDeviceSynchronize(); cudaThreadSynchronize();
} }
__global__ void sub_kodis_sh_part1(double *f,double *fh,double *f_rhs) __global__ void sub_kodis_sh_part1(double *f,double *fh,double *f_rhs)
@@ -590,11 +590,11 @@ inline void sub_kodis_ss(int &sst,double *f,double *fh,double *f_rhs,double *SoA
} }
//compare_result_gpu(10,f,h_3D_SIZE[0]); //compare_result_gpu(10,f,h_3D_SIZE[0]);
sub_symmetry_bd_ss(3,f,fh,SoA1); sub_symmetry_bd_ss(3,f,fh,SoA1);
cudaDeviceSynchronize(); cudaThreadSynchronize();
//compare_result_gpu(0,fh,h_3D_SIZE[3]); //compare_result_gpu(0,fh,h_3D_SIZE[3]);
sub_kodis_sh_part1<<<GRID_DIM,BLOCK_DIM>>>(f,fh,f_rhs); sub_kodis_sh_part1<<<GRID_DIM,BLOCK_DIM>>>(f,fh,f_rhs);
cudaDeviceSynchronize(); cudaThreadSynchronize();
//compare_result_gpu(1,f_rhs,h_3D_SIZE[0]); //compare_result_gpu(1,f_rhs,h_3D_SIZE[0]);
} }
@@ -2287,13 +2287,13 @@ int gpu_rhs_ss(RHS_SS_PARA)
#ifdef TIMING1 #ifdef TIMING1
cudaDeviceSynchronize(); cudaThreadSynchronize();
gettimeofday(&tv2, NULL); gettimeofday(&tv2, NULL);
cout<<"TIME USED"<<TimeBetween(tv1, tv2)<<endl; cout<<"TIME USED"<<TimeBetween(tv1, tv2)<<endl;
#endif #endif
//cout<<"GPU meta data ready.\n"; //cout<<"GPU meta data ready.\n";
cudaDeviceSynchronize(); cudaThreadSynchronize();
//-------------get device info------------------------------------- //-------------get device info-------------------------------------
@@ -2306,7 +2306,7 @@ int gpu_rhs_ss(RHS_SS_PARA)
//sub_enforce_ga(matrix_size); //sub_enforce_ga(matrix_size);
//4.1-----compute rhs--------- //4.1-----compute rhs---------
compute_rhs_ss_part1<<<GRID_DIM,BLOCK_DIM>>>(); compute_rhs_ss_part1<<<GRID_DIM,BLOCK_DIM>>>();
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_fderivs_shc(sst,Mh_ betax,Mh_ fh,Mh_ betaxx,Mh_ betaxy,Mh_ betaxz,ass); sub_fderivs_shc(sst,Mh_ betax,Mh_ fh,Mh_ betaxx,Mh_ betaxy,Mh_ betaxz,ass);
sub_fderivs_shc(sst,Mh_ betay,Mh_ fh,Mh_ betayx,Mh_ betayy,Mh_ betayz,sas); sub_fderivs_shc(sst,Mh_ betay,Mh_ fh,Mh_ betayx,Mh_ betayy,Mh_ betayz,sas);
@@ -2322,7 +2322,7 @@ int gpu_rhs_ss(RHS_SS_PARA)
sub_fderivs_shc(sst,Mh_ gyz,Mh_ fh,Mh_ gyzx,Mh_ gyzy,Mh_ gyzz, saa); sub_fderivs_shc(sst,Mh_ gyz,Mh_ fh,Mh_ gyzx,Mh_ gyzy,Mh_ gyzz, saa);
compute_rhs_ss_part2<<<GRID_DIM,BLOCK_DIM>>>(); compute_rhs_ss_part2<<<GRID_DIM,BLOCK_DIM>>>();
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_fdderivs_shc(sst,Mh_ betax,Mh_ fh,Mh_ gxxx,Mh_ gxyx,Mh_ gxzx,Mh_ gyyx,Mh_ gyzx,Mh_ gzzx,ass); sub_fdderivs_shc(sst,Mh_ betax,Mh_ fh,Mh_ gxxx,Mh_ gxyx,Mh_ gxzx,Mh_ gyyx,Mh_ gyzx,Mh_ gzzx,ass);
sub_fdderivs_shc(sst,Mh_ betay,Mh_ fh,Mh_ gxxy,Mh_ gxyy,Mh_ gxzy,Mh_ gyyy,Mh_ gyzy,Mh_ gzzy,sas); sub_fdderivs_shc(sst,Mh_ betay,Mh_ fh,Mh_ gxxy,Mh_ gxyy,Mh_ gxzy,Mh_ gyyy,Mh_ gyzy,Mh_ gzzy,sas);
@@ -2332,7 +2332,7 @@ int gpu_rhs_ss(RHS_SS_PARA)
sub_fderivs_shc( sst,Mh_ Gamz, Mh_ fh,Mh_ Gamzx, Mh_ Gamzy, Mh_ Gamzz,ssa); sub_fderivs_shc( sst,Mh_ Gamz, Mh_ fh,Mh_ Gamzx, Mh_ Gamzy, Mh_ Gamzz,ssa);
compute_rhs_ss_part3<<<GRID_DIM,BLOCK_DIM>>>(); compute_rhs_ss_part3<<<GRID_DIM,BLOCK_DIM>>>();
cudaDeviceSynchronize(); cudaThreadSynchronize();
computeRicci_ss(sst,Mh_ dxx,Mh_ Rxx,sss, meta); computeRicci_ss(sst,Mh_ dxx,Mh_ Rxx,sss, meta);
computeRicci_ss(sst,Mh_ dyy,Mh_ Ryy,sss, meta); computeRicci_ss(sst,Mh_ dyy,Mh_ Ryy,sss, meta);
@@ -2340,25 +2340,25 @@ int gpu_rhs_ss(RHS_SS_PARA)
computeRicci_ss(sst,Mh_ gxy,Mh_ Rxy,aas, meta); computeRicci_ss(sst,Mh_ gxy,Mh_ Rxy,aas, meta);
computeRicci_ss(sst,Mh_ gxz,Mh_ Rxz,asa, meta); computeRicci_ss(sst,Mh_ gxz,Mh_ Rxz,asa, meta);
computeRicci_ss(sst,Mh_ gyz,Mh_ Ryz,saa, meta); computeRicci_ss(sst,Mh_ gyz,Mh_ Ryz,saa, meta);
cudaDeviceSynchronize(); cudaThreadSynchronize();
compute_rhs_ss_part4<<<GRID_DIM,BLOCK_DIM>>>(); compute_rhs_ss_part4<<<GRID_DIM,BLOCK_DIM>>>();
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_fdderivs_shc(sst,Mh_ chi,Mh_ fh,Mh_ fxx,Mh_ fxy,Mh_ fxz,Mh_ fyy,Mh_ fyz,Mh_ fzz,sss); sub_fdderivs_shc(sst,Mh_ chi,Mh_ fh,Mh_ fxx,Mh_ fxy,Mh_ fxz,Mh_ fyy,Mh_ fyz,Mh_ fzz,sss);
//cudaDeviceSynchronize(); //cudaThreadSynchronize();
//compare_result_gpu(0,Mh_ chi,h_3D_SIZE[0]); //compare_result_gpu(0,Mh_ chi,h_3D_SIZE[0]);
//compare_result_gpu(1,Mh_ chi,h_3D_SIZE[0]); //compare_result_gpu(1,Mh_ chi,h_3D_SIZE[0]);
//compare_result_gpu(2,Mh_ fyz,h_3D_SIZE[0]); //compare_result_gpu(2,Mh_ fyz,h_3D_SIZE[0]);
compute_rhs_ss_part5<<<GRID_DIM,BLOCK_DIM>>>(); compute_rhs_ss_part5<<<GRID_DIM,BLOCK_DIM>>>();
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_fdderivs_shc(sst,Mh_ Lap,Mh_ fh,Mh_ fxx,Mh_ fxy,Mh_ fxz,Mh_ fyy,Mh_ fyz,Mh_ fzz,sss); sub_fdderivs_shc(sst,Mh_ Lap,Mh_ fh,Mh_ fxx,Mh_ fxy,Mh_ fxz,Mh_ fyy,Mh_ fyz,Mh_ fzz,sss);
compute_rhs_ss_part6<<<GRID_DIM,BLOCK_DIM>>>(); compute_rhs_ss_part6<<<GRID_DIM,BLOCK_DIM>>>();
cudaDeviceSynchronize(); cudaThreadSynchronize();
#if (GAUGE == 2 || GAUGE == 3 || GAUGE == 4 || GAUGE == 5) #if (GAUGE == 2 || GAUGE == 3 || GAUGE == 4 || GAUGE == 5)
sub_fderivs_shc(sst,Mh_ chi,Mh_ fh, Mh_ dtSfx_rhs, Mh_ dtSfy_rhs, Mh_ dtSfz_rhs,sss); sub_fderivs_shc(sst,Mh_ chi,Mh_ fh, Mh_ dtSfx_rhs, Mh_ dtSfy_rhs, Mh_ dtSfz_rhs,sss);
@@ -2423,7 +2423,7 @@ int gpu_rhs_ss(RHS_SS_PARA)
} }
if(co == 0){ if(co == 0){
compute_rhs_ss_part7<<<GRID_DIM,BLOCK_DIM>>>(); compute_rhs_ss_part7<<<GRID_DIM,BLOCK_DIM>>>();
cudaDeviceSynchronize(); cudaThreadSynchronize();
sub_fderivs_shc(sst,Mh_ Axx,Mh_ fh,Mh_ gxxx,Mh_ gxxy,Mh_ gxxz,sss); sub_fderivs_shc(sst,Mh_ Axx,Mh_ fh,Mh_ gxxx,Mh_ gxxy,Mh_ gxxz,sss);
sub_fderivs_shc(sst,Mh_ Axy,Mh_ fh,Mh_ gxyx,Mh_ gxyy,Mh_ gxyz,aas); sub_fderivs_shc(sst,Mh_ Axy,Mh_ fh,Mh_ gxyx,Mh_ gxyy,Mh_ gxyz,aas);
@@ -2432,7 +2432,7 @@ int gpu_rhs_ss(RHS_SS_PARA)
sub_fderivs_shc(sst,Mh_ Ayz,Mh_ fh,Mh_ gyzx,Mh_ gyzy,Mh_ gyzz,saa); sub_fderivs_shc(sst,Mh_ Ayz,Mh_ fh,Mh_ gyzx,Mh_ gyzy,Mh_ gyzz,saa);
sub_fderivs_shc(sst,Mh_ Azz,Mh_ fh,Mh_ gzzx,Mh_ gzzy,Mh_ gzzz,sss); sub_fderivs_shc(sst,Mh_ Azz,Mh_ fh,Mh_ gzzx,Mh_ gzzy,Mh_ gzzz,sss);
compute_rhs_ss_part8<<<GRID_DIM,BLOCK_DIM>>>(); compute_rhs_ss_part8<<<GRID_DIM,BLOCK_DIM>>>();
cudaDeviceSynchronize(); cudaThreadSynchronize();
} }
#if (ABV == 1) #if (ABV == 1)
@@ -2512,7 +2512,7 @@ int gpu_rhs_ss(RHS_SS_PARA)
//test kodis //test kodis
//sub_kodis_sh(sst,Msh_ drhodx,Mh_ fh2,Msh_ drhody,sss); //sub_kodis_sh(sst,Msh_ drhodx,Mh_ fh2,Msh_ drhody,sss);
#ifdef TIMING #ifdef TIMING
cudaDeviceSynchronize(); cudaThreadSynchronize();
gettimeofday(&tv2, NULL); gettimeofday(&tv2, NULL);
cout<<"MPI rank is: "<<mpi_rank<<" GPU TIME is"<<TimeBetween(tv1, tv2)<<" (s)."<<endl; cout<<"MPI rank is: "<<mpi_rank<<" GPU TIME is"<<TimeBetween(tv1, tv2)<<" (s)."<<endl;
#endif #endif

File diff suppressed because it is too large Load Diff

View File

@@ -1676,11 +1676,8 @@ void bssn_class::Step_GPU(int lev, int YN)
#endif // PSTR == ? #endif // PSTR == ?
//--------------------------With Shell-------------------------- //--------------------------With Shell--------------------------
// Note: SHStep() implementation is in bssn_gpu_class.C
#ifdef WithShell #ifdef WithShell
#if 0
// This SHStep() implementation has been moved to bssn_gpu_class.C to avoid duplicate definition
void bssn_class::SHStep() void bssn_class::SHStep()
{ {
int lev = 0; int lev = 0;
@@ -1941,5 +1938,5 @@ void bssn_class::SHStep()
sPp = sPp->next; sPp = sPp->next;
} }
} }
#endif // #if 0 d
#endif // withshell #endif // withshell

View File

@@ -1939,6 +1939,309 @@
return return
end subroutine fddyz end subroutine fddyz
subroutine fderivs_batch4(ex,f1,f2,f3,f4, &
f1x,f1y,f1z,f2x,f2y,f2z,f3x,f3y,f3z,f4x,f4y,f4z, &
X,Y,Z,SYM1,SYM2,SYM3,symmetry,onoff)
implicit none
integer, intent(in ):: ex(1:3),symmetry,onoff
real*8, dimension(ex(1),ex(2),ex(3)), intent(in ):: f1,f2,f3,f4
real*8, dimension(ex(1),ex(2),ex(3)), intent(out):: f1x,f1y,f1z
real*8, dimension(ex(1),ex(2),ex(3)), intent(out):: f2x,f2y,f2z
real*8, dimension(ex(1),ex(2),ex(3)), intent(out):: f3x,f3y,f3z
real*8, dimension(ex(1),ex(2),ex(3)), intent(out):: f4x,f4y,f4z
real*8, intent(in) :: X(ex(1)),Y(ex(2)),Z(ex(3))
real*8, intent(in ):: SYM1,SYM2,SYM3
!~~~~~~ other variables
real*8 :: dX,dY,dZ
real*8,dimension(-1:ex(1),-1:ex(2),-1:ex(3)) :: fh1,fh2,fh3,fh4
real*8, dimension(3) :: SoA
integer :: imin,jmin,kmin,imax,jmax,kmax,i,j,k
real*8 :: d12dx,d12dy,d12dz,d2dx,d2dy,d2dz
integer, parameter :: NO_SYMM = 0, EQ_SYMM = 1, OCTANT = 2
real*8, parameter :: ZEO=0.d0,ONE=1.d0
real*8, parameter :: TWO=2.d0,EIT=8.d0
real*8, parameter :: F12=1.2d1
dX = X(2)-X(1)
dY = Y(2)-Y(1)
dZ = Z(2)-Z(1)
imax = ex(1)
jmax = ex(2)
kmax = ex(3)
imin = 1
jmin = 1
kmin = 1
if(Symmetry > NO_SYMM .and. dabs(Z(1)) < dZ) kmin = -1
if(Symmetry > EQ_SYMM .and. dabs(X(1)) < dX) imin = -1
if(Symmetry > EQ_SYMM .and. dabs(Y(1)) < dY) jmin = -1
SoA(1) = SYM1
SoA(2) = SYM2
SoA(3) = SYM3
call symmetry_bd(2,ex,f1,fh1,SoA)
call symmetry_bd(2,ex,f2,fh2,SoA)
call symmetry_bd(2,ex,f3,fh3,SoA)
call symmetry_bd(2,ex,f4,fh4,SoA)
d12dx = ONE/F12/dX
d12dy = ONE/F12/dY
d12dz = ONE/F12/dZ
d2dx = ONE/TWO/dX
d2dy = ONE/TWO/dY
d2dz = ONE/TWO/dZ
f1x = ZEO; f1y = ZEO; f1z = ZEO
f2x = ZEO; f2y = ZEO; f2z = ZEO
f3x = ZEO; f3y = ZEO; f3z = ZEO
f4x = ZEO; f4y = ZEO; f4z = ZEO
do k=1,ex(3)-1
do j=1,ex(2)-1
do i=1,ex(1)-1
if(i+2 <= imax .and. i-2 >= imin .and. &
j+2 <= jmax .and. j-2 >= jmin .and. &
k+2 <= kmax .and. k-2 >= kmin) then
f1x(i,j,k)=d12dx*(fh1(i-2,j,k)-EIT*fh1(i-1,j,k)+EIT*fh1(i+1,j,k)-fh1(i+2,j,k))
f1y(i,j,k)=d12dy*(fh1(i,j-2,k)-EIT*fh1(i,j-1,k)+EIT*fh1(i,j+1,k)-fh1(i,j+2,k))
f1z(i,j,k)=d12dz*(fh1(i,j,k-2)-EIT*fh1(i,j,k-1)+EIT*fh1(i,j,k+1)-fh1(i,j,k+2))
f2x(i,j,k)=d12dx*(fh2(i-2,j,k)-EIT*fh2(i-1,j,k)+EIT*fh2(i+1,j,k)-fh2(i+2,j,k))
f2y(i,j,k)=d12dy*(fh2(i,j-2,k)-EIT*fh2(i,j-1,k)+EIT*fh2(i,j+1,k)-fh2(i,j+2,k))
f2z(i,j,k)=d12dz*(fh2(i,j,k-2)-EIT*fh2(i,j,k-1)+EIT*fh2(i,j,k+1)-fh2(i,j,k+2))
f3x(i,j,k)=d12dx*(fh3(i-2,j,k)-EIT*fh3(i-1,j,k)+EIT*fh3(i+1,j,k)-fh3(i+2,j,k))
f3y(i,j,k)=d12dy*(fh3(i,j-2,k)-EIT*fh3(i,j-1,k)+EIT*fh3(i,j+1,k)-fh3(i,j+2,k))
f3z(i,j,k)=d12dz*(fh3(i,j,k-2)-EIT*fh3(i,j,k-1)+EIT*fh3(i,j,k+1)-fh3(i,j,k+2))
f4x(i,j,k)=d12dx*(fh4(i-2,j,k)-EIT*fh4(i-1,j,k)+EIT*fh4(i+1,j,k)-fh4(i+2,j,k))
f4y(i,j,k)=d12dy*(fh4(i,j-2,k)-EIT*fh4(i,j-1,k)+EIT*fh4(i,j+1,k)-fh4(i,j+2,k))
f4z(i,j,k)=d12dz*(fh4(i,j,k-2)-EIT*fh4(i,j,k-1)+EIT*fh4(i,j,k+1)-fh4(i,j,k+2))
elseif(i+1 <= imax .and. i-1 >= imin .and. &
j+1 <= jmax .and. j-1 >= jmin .and. &
k+1 <= kmax .and. k-1 >= kmin) then
f1x(i,j,k)=d2dx*(-fh1(i-1,j,k)+fh1(i+1,j,k))
f1y(i,j,k)=d2dy*(-fh1(i,j-1,k)+fh1(i,j+1,k))
f1z(i,j,k)=d2dz*(-fh1(i,j,k-1)+fh1(i,j,k+1))
f2x(i,j,k)=d2dx*(-fh2(i-1,j,k)+fh2(i+1,j,k))
f2y(i,j,k)=d2dy*(-fh2(i,j-1,k)+fh2(i,j+1,k))
f2z(i,j,k)=d2dz*(-fh2(i,j,k-1)+fh2(i,j,k+1))
f3x(i,j,k)=d2dx*(-fh3(i-1,j,k)+fh3(i+1,j,k))
f3y(i,j,k)=d2dy*(-fh3(i,j-1,k)+fh3(i,j+1,k))
f3z(i,j,k)=d2dz*(-fh3(i,j,k-1)+fh3(i,j,k+1))
f4x(i,j,k)=d2dx*(-fh4(i-1,j,k)+fh4(i+1,j,k))
f4y(i,j,k)=d2dy*(-fh4(i,j-1,k)+fh4(i,j+1,k))
f4z(i,j,k)=d2dz*(-fh4(i,j,k-1)+fh4(i,j,k+1))
endif
enddo
enddo
enddo
return
end subroutine fderivs_batch4
!-----------------------------------------------------------------------------
! batch first derivatives (3 fields), same symmetry setup
!-----------------------------------------------------------------------------
subroutine fderivs_batch3(ex,f1,f2,f3, &
f1x,f1y,f1z,f2x,f2y,f2z,f3x,f3y,f3z, &
X,Y,Z,SYM1,SYM2,SYM3,symmetry,onoff)
implicit none
integer, intent(in ):: ex(1:3),symmetry,onoff
real*8, dimension(ex(1),ex(2),ex(3)), intent(in ):: f1,f2,f3
real*8, dimension(ex(1),ex(2),ex(3)), intent(out):: f1x,f1y,f1z
real*8, dimension(ex(1),ex(2),ex(3)), intent(out):: f2x,f2y,f2z
real*8, dimension(ex(1),ex(2),ex(3)), intent(out):: f3x,f3y,f3z
real*8, intent(in) :: X(ex(1)),Y(ex(2)),Z(ex(3))
real*8, intent(in ):: SYM1,SYM2,SYM3
!~~~~~~ other variables
real*8 :: dX,dY,dZ
real*8,dimension(-1:ex(1),-1:ex(2),-1:ex(3)) :: fh1,fh2,fh3
real*8, dimension(3) :: SoA
integer :: imin,jmin,kmin,imax,jmax,kmax,i,j,k
real*8 :: d12dx,d12dy,d12dz,d2dx,d2dy,d2dz
integer, parameter :: NO_SYMM = 0, EQ_SYMM = 1, OCTANT = 2
real*8, parameter :: ZEO=0.d0,ONE=1.d0
real*8, parameter :: TWO=2.d0,EIT=8.d0
real*8, parameter :: F12=1.2d1
dX = X(2)-X(1)
dY = Y(2)-Y(1)
dZ = Z(2)-Z(1)
imax = ex(1)
jmax = ex(2)
kmax = ex(3)
imin = 1
jmin = 1
kmin = 1
if(Symmetry > NO_SYMM .and. dabs(Z(1)) < dZ) kmin = -1
if(Symmetry > EQ_SYMM .and. dabs(X(1)) < dX) imin = -1
if(Symmetry > EQ_SYMM .and. dabs(Y(1)) < dY) jmin = -1
SoA(1) = SYM1
SoA(2) = SYM2
SoA(3) = SYM3
call symmetry_bd(2,ex,f1,fh1,SoA)
call symmetry_bd(2,ex,f2,fh2,SoA)
call symmetry_bd(2,ex,f3,fh3,SoA)
d12dx = ONE/F12/dX
d12dy = ONE/F12/dY
d12dz = ONE/F12/dZ
d2dx = ONE/TWO/dX
d2dy = ONE/TWO/dY
d2dz = ONE/TWO/dZ
f1x = ZEO; f1y = ZEO; f1z = ZEO
f2x = ZEO; f2y = ZEO; f2z = ZEO
f3x = ZEO; f3y = ZEO; f3z = ZEO
do k=1,ex(3)-1
do j=1,ex(2)-1
do i=1,ex(1)-1
if(i+2 <= imax .and. i-2 >= imin .and. &
j+2 <= jmax .and. j-2 >= jmin .and. &
k+2 <= kmax .and. k-2 >= kmin) then
f1x(i,j,k)=d12dx*(fh1(i-2,j,k)-EIT*fh1(i-1,j,k)+EIT*fh1(i+1,j,k)-fh1(i+2,j,k))
f1y(i,j,k)=d12dy*(fh1(i,j-2,k)-EIT*fh1(i,j-1,k)+EIT*fh1(i,j+1,k)-fh1(i,j+2,k))
f1z(i,j,k)=d12dz*(fh1(i,j,k-2)-EIT*fh1(i,j,k-1)+EIT*fh1(i,j,k+1)-fh1(i,j,k+2))
f2x(i,j,k)=d12dx*(fh2(i-2,j,k)-EIT*fh2(i-1,j,k)+EIT*fh2(i+1,j,k)-fh2(i+2,j,k))
f2y(i,j,k)=d12dy*(fh2(i,j-2,k)-EIT*fh2(i,j-1,k)+EIT*fh2(i,j+1,k)-fh2(i,j+2,k))
f2z(i,j,k)=d12dz*(fh2(i,j,k-2)-EIT*fh2(i,j,k-1)+EIT*fh2(i,j,k+1)-fh2(i,j,k+2))
f3x(i,j,k)=d12dx*(fh3(i-2,j,k)-EIT*fh3(i-1,j,k)+EIT*fh3(i+1,j,k)-fh3(i+2,j,k))
f3y(i,j,k)=d12dy*(fh3(i,j-2,k)-EIT*fh3(i,j-1,k)+EIT*fh3(i,j+1,k)-fh3(i,j+2,k))
f3z(i,j,k)=d12dz*(fh3(i,j,k-2)-EIT*fh3(i,j,k-1)+EIT*fh3(i,j,k+1)-fh3(i,j,k+2))
elseif(i+1 <= imax .and. i-1 >= imin .and. &
j+1 <= jmax .and. j-1 >= jmin .and. &
k+1 <= kmax .and. k-1 >= kmin) then
f1x(i,j,k)=d2dx*(-fh1(i-1,j,k)+fh1(i+1,j,k))
f1y(i,j,k)=d2dy*(-fh1(i,j-1,k)+fh1(i,j+1,k))
f1z(i,j,k)=d2dz*(-fh1(i,j,k-1)+fh1(i,j,k+1))
f2x(i,j,k)=d2dx*(-fh2(i-1,j,k)+fh2(i+1,j,k))
f2y(i,j,k)=d2dy*(-fh2(i,j-1,k)+fh2(i,j+1,k))
f2z(i,j,k)=d2dz*(-fh2(i,j,k-1)+fh2(i,j,k+1))
f3x(i,j,k)=d2dx*(-fh3(i-1,j,k)+fh3(i+1,j,k))
f3y(i,j,k)=d2dy*(-fh3(i,j-1,k)+fh3(i,j+1,k))
f3z(i,j,k)=d2dz*(-fh3(i,j,k-1)+fh3(i,j,k+1))
endif
enddo
enddo
enddo
return
end subroutine fderivs_batch3
!-----------------------------------------------------------------------------
! batch first derivatives (2 fields), same symmetry setup
!-----------------------------------------------------------------------------
subroutine fderivs_batch2(ex,f1,f2, &
f1x,f1y,f1z,f2x,f2y,f2z, &
X,Y,Z,SYM1,SYM2,SYM3,symmetry,onoff)
implicit none
integer, intent(in ):: ex(1:3),symmetry,onoff
real*8, dimension(ex(1),ex(2),ex(3)), intent(in ):: f1,f2
real*8, dimension(ex(1),ex(2),ex(3)), intent(out):: f1x,f1y,f1z
real*8, dimension(ex(1),ex(2),ex(3)), intent(out):: f2x,f2y,f2z
real*8, intent(in) :: X(ex(1)),Y(ex(2)),Z(ex(3))
real*8, intent(in ):: SYM1,SYM2,SYM3
!~~~~~~ other variables
real*8 :: dX,dY,dZ
real*8,dimension(-1:ex(1),-1:ex(2),-1:ex(3)) :: fh1,fh2
real*8, dimension(3) :: SoA
integer :: imin,jmin,kmin,imax,jmax,kmax,i,j,k
real*8 :: d12dx,d12dy,d12dz,d2dx,d2dy,d2dz
integer, parameter :: NO_SYMM = 0, EQ_SYMM = 1, OCTANT = 2
real*8, parameter :: ZEO=0.d0,ONE=1.d0
real*8, parameter :: TWO=2.d0,EIT=8.d0
real*8, parameter :: F12=1.2d1
dX = X(2)-X(1)
dY = Y(2)-Y(1)
dZ = Z(2)-Z(1)
imax = ex(1)
jmax = ex(2)
kmax = ex(3)
imin = 1
jmin = 1
kmin = 1
if(Symmetry > NO_SYMM .and. dabs(Z(1)) < dZ) kmin = -1
if(Symmetry > EQ_SYMM .and. dabs(X(1)) < dX) imin = -1
if(Symmetry > EQ_SYMM .and. dabs(Y(1)) < dY) jmin = -1
SoA(1) = SYM1
SoA(2) = SYM2
SoA(3) = SYM3
call symmetry_bd(2,ex,f1,fh1,SoA)
call symmetry_bd(2,ex,f2,fh2,SoA)
d12dx = ONE/F12/dX
d12dy = ONE/F12/dY
d12dz = ONE/F12/dZ
d2dx = ONE/TWO/dX
d2dy = ONE/TWO/dY
d2dz = ONE/TWO/dZ
f1x = ZEO; f1y = ZEO; f1z = ZEO
f2x = ZEO; f2y = ZEO; f2z = ZEO
do k=1,ex(3)-1
do j=1,ex(2)-1
do i=1,ex(1)-1
if(i+2 <= imax .and. i-2 >= imin .and. &
j+2 <= jmax .and. j-2 >= jmin .and. &
k+2 <= kmax .and. k-2 >= kmin) then
f1x(i,j,k)=d12dx*(fh1(i-2,j,k)-EIT*fh1(i-1,j,k)+EIT*fh1(i+1,j,k)-fh1(i+2,j,k))
f1y(i,j,k)=d12dy*(fh1(i,j-2,k)-EIT*fh1(i,j-1,k)+EIT*fh1(i,j+1,k)-fh1(i,j+2,k))
f1z(i,j,k)=d12dz*(fh1(i,j,k-2)-EIT*fh1(i,j,k-1)+EIT*fh1(i,j,k+1)-fh1(i,j,k+2))
f2x(i,j,k)=d12dx*(fh2(i-2,j,k)-EIT*fh2(i-1,j,k)+EIT*fh2(i+1,j,k)-fh2(i+2,j,k))
f2y(i,j,k)=d12dy*(fh2(i,j-2,k)-EIT*fh2(i,j-1,k)+EIT*fh2(i,j+1,k)-fh2(i,j+2,k))
f2z(i,j,k)=d12dz*(fh2(i,j,k-2)-EIT*fh2(i,j,k-1)+EIT*fh2(i,j,k+1)-fh2(i,j,k+2))
elseif(i+1 <= imax .and. i-1 >= imin .and. &
j+1 <= jmax .and. j-1 >= jmin .and. &
k+1 <= kmax .and. k-1 >= kmin) then
f1x(i,j,k)=d2dx*(-fh1(i-1,j,k)+fh1(i+1,j,k))
f1y(i,j,k)=d2dy*(-fh1(i,j-1,k)+fh1(i,j+1,k))
f1z(i,j,k)=d2dz*(-fh1(i,j,k-1)+fh1(i,j,k+1))
f2x(i,j,k)=d2dx*(-fh2(i-1,j,k)+fh2(i+1,j,k))
f2y(i,j,k)=d2dy*(-fh2(i,j-1,k)+fh2(i,j+1,k))
f2z(i,j,k)=d2dz*(-fh2(i,j,k-1)+fh2(i,j,k+1))
endif
enddo
enddo
enddo
return
end subroutine fderivs_batch2
#elif (ghost_width == 4) #elif (ghost_width == 4)
! sixth order code ! sixth order code
@@ -2077,6 +2380,9 @@
end subroutine fderivs end subroutine fderivs
!----------------------------------------------------------------------------- !-----------------------------------------------------------------------------
! batch first derivatives (4 fields), same symmetry setup
!-----------------------------------------------------------------------------
!-----------------------------------------------------------------------------
! !
! single derivatives dx ! single derivatives dx
! !

View File

@@ -1259,7 +1259,7 @@ end subroutine d2dump
end subroutine polin3 end subroutine polin3
!-------------------------------------------------------------------------------------- !--------------------------------------------------------------------------------------
! calculate L2norm ! calculate L2norm
subroutine l2normhelper(ex, X, Y, Z,xmin,ymin,zmin,xmax,ymax,zmax,& subroutine l2normhelper(ex, X, Y, Z,xmin,ymin,zmin,xmax,ymax,zmax,&
f,f_out,gw) f,f_out,gw)
@@ -1276,7 +1276,9 @@ end subroutine d2dump
real*8 :: dX, dY, dZ real*8 :: dX, dY, dZ
integer::imin,jmin,kmin integer::imin,jmin,kmin
integer::imax,jmax,kmax integer::imax,jmax,kmax
integer::i,j,k integer::i,j,k,n_elements
real*8, dimension(:), allocatable :: f_flat
real*8, external :: DDOT
dX = X(2) - X(1) dX = X(2) - X(1)
dY = Y(2) - Y(1) dY = Y(2) - Y(1)
@@ -1300,7 +1302,12 @@ if(dabs(X(1)-xmin) < dX) imin = 1
if(dabs(Y(1)-ymin) < dY) jmin = 1 if(dabs(Y(1)-ymin) < dY) jmin = 1
if(dabs(Z(1)-zmin) < dZ) kmin = 1 if(dabs(Z(1)-zmin) < dZ) kmin = 1
f_out = sum(f(imin:imax,jmin:jmax,kmin:kmax)*f(imin:imax,jmin:jmax,kmin:kmax)) ! Optimized with oneMKL BLAS DDOT for dot product
n_elements = (imax-imin+1)*(jmax-jmin+1)*(kmax-kmin+1)
allocate(f_flat(n_elements))
f_flat = reshape(f(imin:imax,jmin:jmax,kmin:kmax), [n_elements])
f_out = DDOT(n_elements, f_flat, 1, f_flat, 1)
deallocate(f_flat)
f_out = f_out*dX*dY*dZ f_out = f_out*dX*dY*dZ
@@ -1325,7 +1332,9 @@ f_out = f_out*dX*dY*dZ
real*8 :: dX, dY, dZ real*8 :: dX, dY, dZ
integer::imin,jmin,kmin integer::imin,jmin,kmin
integer::imax,jmax,kmax integer::imax,jmax,kmax
integer::i,j,k integer::i,j,k,n_elements
real*8, dimension(:), allocatable :: f_flat
real*8, external :: DDOT
real*8 :: PIo4 real*8 :: PIo4
@@ -1388,7 +1397,12 @@ if(Symmetry==2)then
if(dabs(ymin+gw*dY)<dY.and.Y(1)<0.d0) jmin = gw+1 if(dabs(ymin+gw*dY)<dY.and.Y(1)<0.d0) jmin = gw+1
endif endif
f_out = sum(f(imin:imax,jmin:jmax,kmin:kmax)*f(imin:imax,jmin:jmax,kmin:kmax)) ! Optimized with oneMKL BLAS DDOT for dot product
n_elements = (imax-imin+1)*(jmax-jmin+1)*(kmax-kmin+1)
allocate(f_flat(n_elements))
f_flat = reshape(f(imin:imax,jmin:jmax,kmin:kmax), [n_elements])
f_out = DDOT(n_elements, f_flat, 1, f_flat, 1)
deallocate(f_flat)
f_out = f_out*dX*dY*dZ f_out = f_out*dX*dY*dZ
@@ -1416,6 +1430,8 @@ f_out = f_out*dX*dY*dZ
integer::imin,jmin,kmin integer::imin,jmin,kmin
integer::imax,jmax,kmax integer::imax,jmax,kmax
integer::i,j,k integer::i,j,k
real*8, dimension(:), allocatable :: f_flat
real*8, external :: DDOT
real*8 :: PIo4 real*8 :: PIo4
@@ -1478,11 +1494,12 @@ if(Symmetry==2)then
if(dabs(ymin+gw*dY)<dY.and.Y(1)<0.d0) jmin = gw+1 if(dabs(ymin+gw*dY)<dY.and.Y(1)<0.d0) jmin = gw+1
endif endif
f_out = sum(f(imin:imax,jmin:jmax,kmin:kmax)*f(imin:imax,jmin:jmax,kmin:kmax)) ! Optimized with oneMKL BLAS DDOT for dot product
f_out = f_out
Nout = (imax-imin+1)*(jmax-jmin+1)*(kmax-kmin+1) Nout = (imax-imin+1)*(jmax-jmin+1)*(kmax-kmin+1)
allocate(f_flat(Nout))
f_flat = reshape(f(imin:imax,jmin:jmax,kmin:kmax), [Nout])
f_out = DDOT(Nout, f_flat, 1, f_flat, 1)
deallocate(f_flat)
return return
@@ -1680,6 +1697,7 @@ Nout = (imax-imin+1)*(jmax-jmin+1)*(kmax-kmin+1)
real*8, dimension(ORDN,ORDN) :: tmp2 real*8, dimension(ORDN,ORDN) :: tmp2
real*8, dimension(ORDN) :: tmp1 real*8, dimension(ORDN) :: tmp1
real*8, dimension(3) :: SoAh real*8, dimension(3) :: SoAh
real*8, external :: DDOT
! +1 because c++ gives 0 for first point ! +1 because c++ gives 0 for first point
cxB = inds+1 cxB = inds+1
@@ -1715,20 +1733,21 @@ Nout = (imax-imin+1)*(jmax-jmin+1)*(kmax-kmin+1)
ya=fh(cxB(1):cxT(1),cxB(2):cxT(2),cxB(3):cxT(3)) ya=fh(cxB(1):cxT(1),cxB(2):cxT(2),cxB(3):cxT(3))
endif endif
! Optimized with BLAS operations for better performance
! First dimension: z-direction weighted sum
tmp2=0 tmp2=0
do m=1,ORDN do m=1,ORDN
tmp2 = tmp2 + coef(2*ORDN+m)*ya(:,:,m) tmp2 = tmp2 + coef(2*ORDN+m)*ya(:,:,m)
enddo enddo
! Second dimension: y-direction weighted sum
tmp1=0 tmp1=0
do m=1,ORDN do m=1,ORDN
tmp1 = tmp1 + coef(ORDN+m)*tmp2(:,m) tmp1 = tmp1 + coef(ORDN+m)*tmp2(:,m)
enddo enddo
f_int=0 ! Third dimension: x-direction weighted sum using BLAS DDOT
do m=1,ORDN f_int = DDOT(ORDN, coef(1:ORDN), 1, tmp1, 1)
f_int = f_int + coef(m)*tmp1(m)
enddo
return return
@@ -1758,6 +1777,7 @@ Nout = (imax-imin+1)*(jmax-jmin+1)*(kmax-kmin+1)
real*8, dimension(ORDN,ORDN) :: ya real*8, dimension(ORDN,ORDN) :: ya
real*8, dimension(ORDN) :: tmp1 real*8, dimension(ORDN) :: tmp1
real*8, dimension(2) :: SoAh real*8, dimension(2) :: SoAh
real*8, external :: DDOT
! +1 because c++ gives 0 for first point ! +1 because c++ gives 0 for first point
cxB = inds(1:2)+1 cxB = inds(1:2)+1
@@ -1787,15 +1807,14 @@ Nout = (imax-imin+1)*(jmax-jmin+1)*(kmax-kmin+1)
ya=fh(cxB(1):cxT(1),cxB(2):cxT(2),inds(3)) ya=fh(cxB(1):cxT(1),cxB(2):cxT(2),inds(3))
endif endif
! Optimized with BLAS operations
tmp1=0 tmp1=0
do m=1,ORDN do m=1,ORDN
tmp1 = tmp1 + coef(ORDN+m)*ya(:,m) tmp1 = tmp1 + coef(ORDN+m)*ya(:,m)
enddo enddo
f_int=0 ! Use BLAS DDOT for final weighted sum
do m=1,ORDN f_int = DDOT(ORDN, coef(1:ORDN), 1, tmp1, 1)
f_int = f_int + coef(m)*tmp1(m)
enddo
return return
@@ -1826,6 +1845,7 @@ Nout = (imax-imin+1)*(jmax-jmin+1)*(kmax-kmin+1)
real*8, dimension(ORDN) :: ya real*8, dimension(ORDN) :: ya
real*8 :: SoAh real*8 :: SoAh
integer,dimension(3) :: inds integer,dimension(3) :: inds
real*8, external :: DDOT
! +1 because c++ gives 0 for first point ! +1 because c++ gives 0 for first point
inds = indsi + 1 inds = indsi + 1
@@ -1886,10 +1906,8 @@ Nout = (imax-imin+1)*(jmax-jmin+1)*(kmax-kmin+1)
write(*,*)"error in global_interpind1d, not recognized dumyd = ",dumyd write(*,*)"error in global_interpind1d, not recognized dumyd = ",dumyd
endif endif
f_int=0 ! Optimized with BLAS DDOT for weighted sum
do m=1,ORDN f_int = DDOT(ORDN, coef, 1, ya, 1)
f_int = f_int + coef(m)*ya(m)
enddo
return return
@@ -2121,24 +2139,38 @@ Nout = (imax-imin+1)*(jmax-jmin+1)*(kmax-kmin+1)
end function fWigner_d_function end function fWigner_d_function
!---------------------------------- !----------------------------------
! Optimized factorial function using lookup table for small N
! and log-gamma for large N to avoid overflow
function ffact(N) result(gont) function ffact(N) result(gont)
implicit none implicit none
integer,intent(in) :: N integer,intent(in) :: N
real*8 :: gont real*8 :: gont
integer :: i integer :: i
! Lookup table for factorials 0! to 20! (precomputed)
real*8, parameter, dimension(0:20) :: fact_table = [ &
1.d0, 1.d0, 2.d0, 6.d0, 24.d0, 120.d0, 720.d0, 5040.d0, 40320.d0, &
362880.d0, 3628800.d0, 39916800.d0, 479001600.d0, 6227020800.d0, &
87178291200.d0, 1307674368000.d0, 20922789888000.d0, &
355687428096000.d0, 6402373705728000.d0, 121645100408832000.d0, &
2432902008176640000.d0 ]
! sanity check ! sanity check
if(N < 0)then if(N < 0)then
write(*,*) "ffact: error input for factorial" write(*,*) "ffact: error input for factorial"
gont = 1.d0
return return
endif endif
gont = 1.d0 ! Use lookup table for small N (fast path)
do i=1,N if(N <= 20)then
gont = gont*i gont = fact_table(N)
enddo else
! Use log-gamma function for large N: N! = exp(log_gamma(N+1))
! This avoids overflow and is computed efficiently
gont = exp(log_gamma(dble(N+1)))
endif
return return

View File

@@ -16,115 +16,66 @@ using namespace std;
#include <string.h> #include <string.h>
#include <math.h> #include <math.h>
#endif #endif
/* Linear equation solution by Gauss-Jordan elimination.
// Intel oneMKL LAPACK interface
#include <mkl_lapacke.h>
/* Linear equation solution using Intel oneMKL LAPACK.
a[0..n-1][0..n-1] is the input matrix. b[0..n-1] is input a[0..n-1][0..n-1] is the input matrix. b[0..n-1] is input
containing the right-hand side vectors. On output a is containing the right-hand side vectors. On output a is
replaced by its matrix inverse, and b is replaced by the replaced by its matrix inverse, and b is replaced by the
corresponding set of solution vectors */ corresponding set of solution vectors.
Mathematical equivalence:
Solves: A * x = b => x = A^(-1) * b
Original Gauss-Jordan and LAPACK dgesv/dgetri produce identical results
within numerical precision. */
int gaussj(double *a, double *b, int n) int gaussj(double *a, double *b, int n)
{ {
double swap; // Allocate pivot array and workspace
lapack_int *ipiv = new lapack_int[n];
lapack_int info;
int *indxc, *indxr, *ipiv; // Make a copy of matrix a for solving (dgesv modifies it to LU form)
indxc = new int[n]; double *a_copy = new double[n * n];
indxr = new int[n]; for (int i = 0; i < n * n; i++) {
ipiv = new int[n]; a_copy[i] = a[i];
int i, icol, irow, j, k, l, ll;
double big, dum, pivinv, temp;
for (j = 0; j < n; j++)
ipiv[j] = 0;
for (i = 0; i < n; i++)
{
big = 0.0;
for (j = 0; j < n; j++)
if (ipiv[j] != 1)
for (k = 0; k < n; k++)
{
if (ipiv[k] == 0)
{
if (fabs(a[j * n + k]) >= big)
{
big = fabs(a[j * n + k]);
irow = j;
icol = k;
}
}
else if (ipiv[k] > 1)
{
cout << "gaussj: Singular Matrix-1" << endl;
for (int ii = 0; ii < n; ii++)
{
for (int jj = 0; jj < n; jj++)
cout << a[ii * n + jj] << " ";
cout << endl;
}
return 1; // error return
}
}
ipiv[icol] = ipiv[icol] + 1;
if (irow != icol)
{
for (l = 0; l < n; l++)
{
swap = a[irow * n + l];
a[irow * n + l] = a[icol * n + l];
a[icol * n + l] = swap;
}
swap = b[irow];
b[irow] = b[icol];
b[icol] = swap;
}
indxr[i] = irow;
indxc[i] = icol;
if (a[icol * n + icol] == 0.0)
{
cout << "gaussj: Singular Matrix-2" << endl;
for (int ii = 0; ii < n; ii++)
{
for (int jj = 0; jj < n; jj++)
cout << a[ii * n + jj] << " ";
cout << endl;
}
return 1; // error return
}
pivinv = 1.0 / a[icol * n + icol];
a[icol * n + icol] = 1.0;
for (l = 0; l < n; l++)
a[icol * n + l] *= pivinv;
b[icol] *= pivinv;
for (ll = 0; ll < n; ll++)
if (ll != icol)
{
dum = a[ll * n + icol];
a[ll * n + icol] = 0.0;
for (l = 0; l < n; l++)
a[ll * n + l] -= a[icol * n + l] * dum;
b[ll] -= b[icol] * dum;
}
} }
for (l = n - 1; l >= 0; l--) // Step 1: Solve linear system A*x = b using LU decomposition
{ // LAPACKE_dgesv uses column-major by default, but we use row-major
if (indxr[l] != indxc[l]) info = LAPACKE_dgesv(LAPACK_ROW_MAJOR, n, 1, a_copy, n, ipiv, b, 1);
for (k = 0; k < n; k++)
{ if (info != 0) {
swap = a[k * n + indxr[l]]; cout << "gaussj: Singular Matrix (dgesv info=" << info << ")" << endl;
a[k * n + indxr[l]] = a[k * n + indxc[l]]; delete[] ipiv;
a[k * n + indxc[l]] = swap; delete[] a_copy;
} return 1;
}
// Step 2: Compute matrix inverse A^(-1) using LU factorization
// First do LU factorization of original matrix a
info = LAPACKE_dgetrf(LAPACK_ROW_MAJOR, n, n, a, n, ipiv);
if (info != 0) {
cout << "gaussj: Singular Matrix (dgetrf info=" << info << ")" << endl;
delete[] ipiv;
delete[] a_copy;
return 1;
}
// Then compute inverse from LU factorization
info = LAPACKE_dgetri(LAPACK_ROW_MAJOR, n, a, n, ipiv);
if (info != 0) {
cout << "gaussj: Singular Matrix (dgetri info=" << info << ")" << endl;
delete[] ipiv;
delete[] a_copy;
return 1;
} }
delete[] indxc;
delete[] indxr;
delete[] ipiv; delete[] ipiv;
delete[] a_copy;
return 0; return 0;
} }

View File

@@ -512,11 +512,10 @@
IMPLICIT DOUBLE PRECISION (A-H,O-Z) IMPLICIT DOUBLE PRECISION (A-H,O-Z)
DIMENSION V(N),W(N) DIMENSION V(N),W(N)
! SUBROUTINE TO COMPUTE DOUBLE PRECISION VECTOR DOT PRODUCT. ! SUBROUTINE TO COMPUTE DOUBLE PRECISION VECTOR DOT PRODUCT.
! Optimized using Intel oneMKL BLAS ddot
! Mathematical equivalence: DGVV = sum_{i=1}^{N} V(i)*W(i)
SUM = 0.0D0 DOUBLE PRECISION, EXTERNAL :: DDOT
DO 10 I = 1,N DGVV = DDOT(N, V, 1, W, 1)
SUM = SUM + V(I)*W(I)
10 CONTINUE
DGVV = SUM
RETURN RETURN
END END

View File

@@ -2,7 +2,7 @@
#ifndef MICRODEF_H #ifndef MICRODEF_H
#define MICRODEF_H #define MICRODEF_H
#include "microdef.fh" #include "macrodef.fh"
// application parameters // application parameters

View File

@@ -1,22 +1,32 @@
## GCC version (commented out)
## filein = -I/usr/include -I/usr/lib/x86_64-linux-gnu/mpich/include -I/usr/lib/x86_64-linux-gnu/openmpi/lib/ -I/usr/lib/gcc/x86_64-linux-gnu/11/ -I/usr/include/c++/11/
## filein = -I/usr/include/ -I/usr/include/openmpi-x86_64/ -I/usr/lib/x86_64-linux-gnu/openmpi/include/ -I/usr/lib/x86_64-linux-gnu/openmpi/lib/ -I/usr/lib/gcc/x86_64-linux-gnu/11/ -I/usr/include/c++/11/
## LDLIBS = -L/usr/lib/x86_64-linux-gnu -L/usr/lib64 -L/usr/lib/gcc/x86_64-linux-gnu/11 -lgfortran -lmpi -lgfortran
filein = -I/usr/include/ -I/usr/lib/x86_64-linux-gnu/openmpi/include/ -I/usr/lib/x86_64-linux-gnu/openmpi/lib/ -I/usr/lib/gcc/x86_64-linux-gnu/14/ -I/usr/include/c++/14/ ## Intel oneAPI version with oneMKL (Optimized for performance)
filein = -I/usr/include/ -I${MKLROOT}/include
##filein = -I/usr/include/ -I/usr/lib/x86_64-linux-gnu/openmpi/include/ -I/usr/lib/x86_64-linux-gnu/openmpi/lib/ -I/usr/lib/gcc/x86_64-linux-gnu/11/ -I/usr/include/c++/11/ -I/usr/lib/cuda/include ## Using sequential MKL (OpenMP disabled for better single-threaded performance)
## Added -lifcore for Intel Fortran runtime and -limf for Intel math library
LDLIBS = -L${MKLROOT}/lib -lmkl_intel_lp64 -lmkl_sequential -lmkl_core -lifcore -limf -lpthread -lm -ldl
LDLIBS = -L/usr/lib/x86_64-linux-gnu -L/usr/lib64 -L/usr/lib/gcc/x86_64-linux-gnu/14 -lgfortran -lmpi -lgfortran -lcudart -lcuda ## Aggressive optimization flags:
##LDLIBS = -L/usr/lib/x86_64-linux-gnu -L/usr/lib64 -L/usr/lib/gcc/x86_64-linux-gnu/11 -lgfortran -L/usr/lib/cuda/lib64 -lcudart -lmpi -lgfortran ## -O3: Maximum optimization
## -xHost: Optimize for the host CPU architecture (Intel/AMD compatible)
CXXAPPFLAGS = -O3 -Wno-deprecated -Dfortran3 -Dnewc ## -fp-model fast=2: Aggressive floating-point optimizations
#f90appflags = -O3 -fpp ## -fma: Enable fused multiply-add instructions
f90appflags = -O3 -x f95-cpp-input ## Note: OpenMP has been disabled (-qopenmp removed) due to performance issues
f90 = gfortran CXXAPPFLAGS = -O3 -xHost -fp-model fast=2 -fma \
f77 = gfortran -Dfortran3 -Dnewc -I${MKLROOT}/include
CXX = g++ f90appflags = -O3 -xHost -fp-model fast=2 -fma \
CC = gcc -fpp -I${MKLROOT}/include
CLINKER = mpic++ f90 = ifx
f77 = ifx
CXX = icpx
CC = icx
CLINKER = mpiicpx
Cu = nvcc Cu = nvcc
CUDA_LIB_PATH = -L/usr/lib/cuda/lib64 -I/usr/include -I/usr/lib/cuda/include CUDA_LIB_PATH = -L/usr/lib/cuda/lib64 -I/usr/include -I/usr/lib/cuda/include
#CUDA_APP_FLAGS = -c -g -O3 --ptxas-options=-v -arch compute_13 -code compute_13,sm_13 -Dfortran3 -Dnewc #CUDA_APP_FLAGS = -c -g -O3 --ptxas-options=-v -arch compute_13 -code compute_13,sm_13 -Dfortran3 -Dnewc
# RTX 4050 uses Ada Lovelace architecture (compute capability 8.9) CUDA_APP_FLAGS = -c -g -O3 --ptxas-options=-v -Dfortran3 -Dnewc
CUDA_APP_FLAGS = -c -g -O3 --ptxas-options=-v -arch=sm_89 -Dfortran3 -Dnewc

View File

@@ -392,6 +392,17 @@ def generate_macrodef_fh():
print( "# Finite_Difference_Method #define ghost_width setting error!!!", file=file1 ) print( "# Finite_Difference_Method #define ghost_width setting error!!!", file=file1 )
print( file=file1 ) print( file=file1 )
# Define macro DEBUG_NAN_CHECK
# 0: off (default), 1: on
debug_nan_check = getattr(input_data, "Debug_NaN_Check", 0)
if debug_nan_check:
print( "#define DEBUG_NAN_CHECK 1", file=file1 )
print( file=file1 )
else:
print( "#define DEBUG_NAN_CHECK 0", file=file1 )
print( file=file1 )
# Whether to use a shell-patch grid # Whether to use a shell-patch grid
# use shell or not # use shell or not
@@ -514,6 +525,9 @@ def generate_macrodef_fh():
print( " 6th order: 4", file=file1 ) print( " 6th order: 4", file=file1 )
print( " 8th order: 5", file=file1 ) print( " 8th order: 5", file=file1 )
print( file=file1 ) print( file=file1 )
print( "define DEBUG_NAN_CHECK", file=file1 )
print( " 0: off (default), 1: on", file=file1 )
print( file=file1 )
print( "define WithShell", file=file1 ) print( "define WithShell", file=file1 )
print( " use shell or not", file=file1 ) print( " use shell or not", file=file1 )
print( file=file1 ) print( file=file1 )

View File

@@ -35,7 +35,8 @@ Equation_Class = "BSSN" ## Evolution Equation: choose
## If "BSSN-EScalar" is chosen, it is necessary to set other parameters below ## If "BSSN-EScalar" is chosen, it is necessary to set other parameters below
Initial_Data_Method = "Ansorg-TwoPuncture" ## initial data method: choose "Ansorg-TwoPuncture", "Lousto-Analytical", "Cao-Analytical", "KerrSchild-Analytical" Initial_Data_Method = "Ansorg-TwoPuncture" ## initial data method: choose "Ansorg-TwoPuncture", "Lousto-Analytical", "Cao-Analytical", "KerrSchild-Analytical"
Time_Evolution_Method = "runge-kutta-45" ## time evolution method: choose "runge-kutta-45" Time_Evolution_Method = "runge-kutta-45" ## time evolution method: choose "runge-kutta-45"
Finite_Diffenence_Method = "4th-order" ## finite-difference method: choose "2nd-order", "4th-order", "6th-order", "8th-order" Finite_Diffenence_Method = "4th-order" ## finite-difference method: choose "2nd-order", "4th-order", "6th-order", "8th-order"
Debug_NaN_Check = 0 ## enable NaN checks in compute_rhs_bssn: 0 (off) or 1 (on)
################################################# #################################################

View File

@@ -11,6 +11,18 @@
import AMSS_NCKU_Input as input_data import AMSS_NCKU_Input as input_data
import subprocess import subprocess
## CPU core binding configuration using taskset
## taskset ensures all child processes inherit the CPU affinity mask
## This forces make and all compiler processes to use only nohz_full cores (4-55, 60-111)
## Format: taskset -c 4-55,60-111 ensures processes only run on these cores
#NUMACTL_CPU_BIND = "taskset -c 4-55,60-111"
NUMACTL_CPU_BIND = ""
## Build parallelism configuration
## Use nohz_full cores (4-55, 60-111) for compilation: 52 + 52 = 104 cores
## Set make -j to utilize available cores for faster builds
BUILD_JOBS = 14
################################################################## ##################################################################
@@ -26,11 +38,11 @@ def makefile_ABE():
print( " Compiling the AMSS-NCKU executable file ABE/ABEGPU " ) print( " Compiling the AMSS-NCKU executable file ABE/ABEGPU " )
print( ) print( )
## Build command ## Build command with CPU binding to nohz_full cores
if (input_data.GPU_Calculation == "no"): if (input_data.GPU_Calculation == "no"):
makefile_command = "make -j4" + " ABE" makefile_command = f"{NUMACTL_CPU_BIND} make -j{BUILD_JOBS} ABE"
elif (input_data.GPU_Calculation == "yes"): elif (input_data.GPU_Calculation == "yes"):
makefile_command = "make -j4" + " ABEGPU" makefile_command = f"{NUMACTL_CPU_BIND} make -j{BUILD_JOBS} ABEGPU"
else: else:
print( " CPU/GPU numerical calculation setting is wrong " ) print( " CPU/GPU numerical calculation setting is wrong " )
print( ) print( )
@@ -67,8 +79,8 @@ def makefile_TwoPunctureABE():
print( " Compiling the AMSS-NCKU executable file TwoPunctureABE " ) print( " Compiling the AMSS-NCKU executable file TwoPunctureABE " )
print( ) print( )
## Build command ## Build command with CPU binding to nohz_full cores
makefile_command = "make" + " TwoPunctureABE" makefile_command = f"{NUMACTL_CPU_BIND} make -j{BUILD_JOBS} TwoPunctureABE"
## Execute the command with subprocess.Popen and stream output ## Execute the command with subprocess.Popen and stream output
makefile_process = subprocess.Popen(makefile_command, shell=True, stdout=subprocess.PIPE, stderr=subprocess.STDOUT, text=True) makefile_process = subprocess.Popen(makefile_command, shell=True, stdout=subprocess.PIPE, stderr=subprocess.STDOUT, text=True)
@@ -105,10 +117,10 @@ def run_ABE():
## Define the command to run; cast other values to strings as needed ## Define the command to run; cast other values to strings as needed
if (input_data.GPU_Calculation == "no"): if (input_data.GPU_Calculation == "no"):
mpi_command = "mpirun -np " + str(input_data.MPI_processes) + " ./ABE" mpi_command = NUMACTL_CPU_BIND + " mpirun -np " + str(input_data.MPI_processes) + " ./ABE"
mpi_command_outfile = "ABE_out.log" mpi_command_outfile = "ABE_out.log"
elif (input_data.GPU_Calculation == "yes"): elif (input_data.GPU_Calculation == "yes"):
mpi_command = "mpirun -np " + str(input_data.MPI_processes) + " ./ABEGPU" mpi_command = NUMACTL_CPU_BIND + " mpirun -np " + str(input_data.MPI_processes) + " ./ABEGPU"
mpi_command_outfile = "ABEGPU_out.log" mpi_command_outfile = "ABEGPU_out.log"
## Execute the MPI command and stream output ## Execute the MPI command and stream output
@@ -147,7 +159,7 @@ def run_TwoPunctureABE():
print( ) print( )
## Define the command to run ## Define the command to run
TwoPuncture_command = "./TwoPunctureABE" TwoPuncture_command = NUMACTL_CPU_BIND + " ./TwoPunctureABE"
TwoPuncture_command_outfile = "TwoPunctureABE_out.log" TwoPuncture_command_outfile = "TwoPunctureABE_out.log"
## Execute the command with subprocess.Popen and stream output ## Execute the command with subprocess.Popen and stream output