diff --git a/components/homme/CMakeLists.txt b/components/homme/CMakeLists.txt index 80a89a296910..6fe81180ab54 100644 --- a/components/homme/CMakeLists.txt +++ b/components/homme/CMakeLists.txt @@ -206,7 +206,9 @@ IF (HOMME_USE_KOKKOS) STRING (TOUPPER ${HOMMEXX_EXEC_SPACE} HOMMEXX_EXEC_SPACE_UPPER) - IF (HOMMEXX_EXEC_SPACE_UPPER STREQUAL "HIP") + IF (${HOMMEXX_EXEC_SPACE_UPPER} STREQUAL "SYCL") + SET (HOMMEXX_SYCL_SPACE ON) + ELSEIF (${HOMMEXX_EXEC_SPACE_UPPER} STREQUAL "HIP") SET (HOMMEXX_HIP_SPACE ON) ELSEIF (HOMMEXX_EXEC_SPACE_UPPER STREQUAL "CUDA") SET (HOMMEXX_CUDA_SPACE ON) @@ -303,7 +305,7 @@ SET (HOMMEXX_ENABLE_GPU_F90 FALSE) IF (HOMME_USE_KOKKOS) - IF (CUDA_BUILD OR HIP_BUILD) + IF (CUDA_BUILD OR HIP_BUILD OR SYCL_BUILD) SET (DEFAULT_VECTOR_SIZE 1) SET (HOMMEXX_ENABLE_GPU TRUE) SET (HOMMEXX_ENABLE_GPU_F90 TRUE) @@ -312,7 +314,7 @@ IF (HOMME_USE_KOKKOS) ENDIF() SET (HOMMEXX_VECTOR_SIZE ${DEFAULT_VECTOR_SIZE} CACHE STRING - "If AVX or Cuda or HIP don't take priority, use this software vector size.") + "If AVX or Cuda or HIP or SYCL don't take priority, use this software vector size.") IF (CMAKE_BUILD_TYPE_UPPER MATCHES "DEBUG" OR CMAKE_BUILD_TYPE_UPPER MATCHES "RELWITHDEBINFO") SET (HOMMEXX_DEBUG ON) diff --git a/components/homme/cmake/machineFiles/aurora-aot.cmake b/components/homme/cmake/machineFiles/aurora-aot.cmake new file mode 100644 index 000000000000..094ec8882784 --- /dev/null +++ b/components/homme/cmake/machineFiles/aurora-aot.cmake @@ -0,0 +1,64 @@ +#module restore +#module load oneapi/eng-compiler/2022.12.30.005 +#module load intel_compute_runtime/release/agama-devel-627 +#module load spack cmake +#module list + + +SET (SUNSPOT_MACHINE TRUE CACHE BOOL "") + +SET(BUILD_HOMME_WITHOUT_PIOLIBRARY TRUE CACHE BOOL "") +SET(HOMMEXX_MPI_ON_DEVICE FALSE CACHE BOOL "") + +SET(HOMME_FIND_BLASLAPACK TRUE CACHE BOOL "") + +SET(WITH_PNETCDF FALSE CACHE FILEPATH "") + +SET(USE_QUEUING FALSE CACHE BOOL "") + +#temp hack +SET(HOMME_USE_KOKKOS TRUE CACHE BOOL "") + +SET(BUILD_HOMME_PREQX_KOKKOS TRUE CACHE BOOL "") +SET(BUILD_HOMME_THETA_KOKKOS TRUE CACHE BOOL "") + +#set(KOKKOS_HOME "/home/onguba/kokkos-build/mar05-aot/install" CACHE STRING "") +#set(E3SM_KOKKOS_PATH ${KOKKOS_HOME} CACHE STRING "") + +SET(USE_TRILINOS OFF CACHE BOOL "") + +SET(SYCL_BUILD TRUE CACHE BOOL "") +SET(HOMME_ENABLE_COMPOSE FALSE CACHE BOOL "") + +SET(CMAKE_CXX_STANDARD 17) + +SET(CMAKE_C_COMPILER "mpicc" CACHE STRING "") +SET(CMAKE_Fortran_COMPILER "mpifort" CACHE STRING "") +SET(CMAKE_CXX_COMPILER "mpicxx" CACHE STRING "") + +# -fsycl-link-huge-device-code for theta to get build +#JIT flags +#SET(SYCL_COMPILE_FLAGS "-std=c++17 -fsycl -fsycl-device-code-split=per_kernel -fno-sycl-id-queries-fit-in-int -fsycl-unnamed-lambda") +#SET(SYCL_LINK_FLAGS "-fsycl -fsycl-link-huge-device-code -fsycl-device-code-split=per_kernel -fsycl-targets=spir64") + +#AOT flags +SET(SYCL_COMPILE_FLAGS "-std=c++17 -fsycl -fsycl-device-code-split=per_kernel -fno-sycl-id-queries-fit-in-int -fsycl-unnamed-lambda") +SET(SYCL_LINK_FLAGS "-fsycl-max-parallel-link-jobs=32 -fsycl-link-huge-device-code -fsycl -fsycl-device-code-split=per_kernel -fsycl-targets=spir64_gen -Xsycl-target-backend \"-device 12.60.7\"") + +SET(ADD_Fortran_FLAGS "-fc=ifx -fpscomp logicals -O3 -DNDEBUG -DCPRINTEL -g" CACHE STRING "") +SET(ADD_C_FLAGS "-O3 -DNDEBUG " CACHE STRING "") + +SET(ADD_CXX_FLAGS "-std=c++17 -O3 -DNDEBUG ${SYCL_COMPILE_FLAGS}" CACHE STRING "") +SET(ADD_LINKER_FLAGS "-O3 -DNDEBUG ${SYCL_LINK_FLAGS} -fortlib" CACHE STRING "") + +set (ENABLE_OPENMP OFF CACHE BOOL "") +set (ENABLE_COLUMN_OPENMP OFF CACHE BOOL "") +set (ENABLE_HORIZ_OPENMP OFF CACHE BOOL "") + +set (HOMME_TESTING_PROFILE "dev" CACHE STRING "") + +set (USE_NUM_PROCS 4 CACHE STRING "") + +SET (USE_MPI_OPTIONS "--bind-to core" CACHE FILEPATH "") + + diff --git a/components/homme/cmake/machineFiles/aurora-jit.cmake b/components/homme/cmake/machineFiles/aurora-jit.cmake new file mode 100644 index 000000000000..1941fa9eb3f3 --- /dev/null +++ b/components/homme/cmake/machineFiles/aurora-jit.cmake @@ -0,0 +1,58 @@ +#module restore +#module load oneapi/eng-compiler/2022.12.30.005 +#module load intel_compute_runtime/release/agama-devel-627 +#module load spack cmake +#module list + + + +SET(BUILD_HOMME_WITHOUT_PIOLIBRARY TRUE CACHE BOOL "") +SET(HOMMEXX_MPI_ON_DEVICE FALSE CACHE BOOL "") + +SET(HOMME_FIND_BLASLAPACK TRUE CACHE BOOL "") + +SET(WITH_PNETCDF FALSE CACHE FILEPATH "") + +SET(USE_QUEUING FALSE CACHE BOOL "") + +#temp hack +SET(HOMME_USE_KOKKOS TRUE CACHE BOOL "") + +SET(BUILD_HOMME_PREQX_KOKKOS TRUE CACHE BOOL "") +SET(BUILD_HOMME_THETA_KOKKOS TRUE CACHE BOOL "") + +#set(KOKKOS_HOME "/home/onguba/kokkos-build/jan03-2024/install" CACHE STRING "") +#set(E3SM_KOKKOS_PATH ${KOKKOS_HOME} CACHE STRING "") + +SET(USE_TRILINOS OFF CACHE BOOL "") + +SET(SYCL_BUILD TRUE CACHE BOOL "") +SET(HOMME_ENABLE_COMPOSE FALSE CACHE BOOL "") + +SET(CMAKE_CXX_STANDARD 17) + +SET(CMAKE_C_COMPILER "mpicc" CACHE STRING "") +SET(CMAKE_Fortran_COMPILER "mpifort" CACHE STRING "") +SET(CMAKE_CXX_COMPILER "mpicxx" CACHE STRING "") + +# -fsycl-link-huge-device-code for theta to get build +SET(SYCL_COMPILE_FLAGS "-std=c++17 -fsycl -fsycl-device-code-split=per_kernel -fno-sycl-id-queries-fit-in-int -fsycl-unnamed-lambda") +SET(SYCL_LINK_FLAGS "-fsycl -fsycl-link-huge-device-code -fsycl-device-code-split=per_kernel -fsycl-targets=spir64") + +SET(ADD_Fortran_FLAGS "-fc=ifx -O3 -DNDEBUG -DCPRINTEL -g" CACHE STRING "") +SET(ADD_C_FLAGS "-O3 -DNDEBUG " CACHE STRING "") + +SET(ADD_CXX_FLAGS "-std=c++17 -O3 -DNDEBUG ${SYCL_COMPILE_FLAGS}" CACHE STRING "") +SET(ADD_LINKER_FLAGS "-O3 -DNDEBUG ${SYCL_LINK_FLAGS} -fortlib" CACHE STRING "") + +set (ENABLE_OPENMP OFF CACHE BOOL "") +set (ENABLE_COLUMN_OPENMP OFF CACHE BOOL "") +set (ENABLE_HORIZ_OPENMP OFF CACHE BOOL "") + +set (HOMME_TESTING_PROFILE "dev" CACHE STRING "") + +set (USE_NUM_PROCS 4 CACHE STRING "") + +SET (USE_MPI_OPTIONS "--bind-to core" CACHE FILEPATH "") + + diff --git a/components/homme/cmake/machineFiles/polaris-a100.sh b/components/homme/cmake/machineFiles/polaris-a100.sh new file mode 100644 index 000000000000..2b63c61a55e7 --- /dev/null +++ b/components/homme/cmake/machineFiles/polaris-a100.sh @@ -0,0 +1,74 @@ +#Currently Loaded Modules: +# 1) craype-x86-rome 6) craype/2.7.15 11) cray-libpals/1.1.7 16) nvhpc-mixed/21.9 +# 2) libfabric/1.11.0.4.125 7) cray-dsmml/0.2.2 12) PrgEnv-gnu/8.3.3 17) cudatoolkit-standalone/11.6.2 +# 3) craype-network-ofi 8) cray-pmi/6.1.2 13) gnu-parallel/2021-09-22 18) cmake/3.23.2 +# 4) perftools-base/22.05.0 9) cray-pmi-lib/6.0.17 14) gcc/11.2.0 +# 5) craype-accel-nvidia80 10) cray-pals/1.1.7 15) cray-mpich/8.1.16 + + + +#SET(HOMMEXX_EXEC_SPACE CUDA CACHE STRING "") +#SET(HOMMEXX_MPI_ON_DEVICE FALSE CACHE BOOL "") +#SET(HOMMEXX_CUDA_MAX_WARP_PER_TEAM "16" CACHE STRING "") + +# cray-hdf5-parallel/1.12.0.6 cray-netcdf-hdf5parallel/4.7.4.6 cray-parallel-netcdf/1.12.1.6 +#SET(NETCDF_DIR $ENV{CRAY_NETCDF_HDF5PARALLEL_PREFIX} CACHE FILEPATH "") +#SET(PNETCDF_DIR $ENV{CRAY_PARALLEL_NETCDF_DIR} CACHE FILEPATH "") +#SET(HDF5_DIR $ENV{CRAY_HDF5_PARALLEL_PREFIX} CACHE FILEPATH "") + +#for scorpio +#SET (NetCDF_C_PATH $ENV{CRAY_NETCDF_HDF5PARALLEL_PREFIX} CACHE FILEPATH "") +#SET (NetCDF_Fortran_PATH $ENV{CRAY_NETCDF_HDF5PARALLEL_PREFIX} CACHE FILEPATH "") + +SET(BUILD_HOMME_WITHOUT_PIOLIBRARY TRUE CACHE BOOL "") + +SET(HOMME_FIND_BLASLAPACK FALSE CACHE BOOL "") + +SET(WITH_PNETCDF FALSE CACHE FILEPATH "") + +SET(USE_QUEUING FALSE CACHE BOOL "") + +SET(BUILD_HOMME_THETA_KOKKOS TRUE CACHE BOOL "") + +SET(CUDA_BUILD TRUE CACHE BOOL "") + +#SET(HOMMEXX_BFB_TESTING TRUE CACHE BOOL "") + +SET(USE_TRILINOS OFF CACHE BOOL "") + +SET(Kokkos_ENABLE_OPENMP OFF CACHE BOOL "") +SET(Kokkos_ENABLE_CUDA ON CACHE BOOL "") +SET(Kokkos_ENABLE_CUDA_LAMBDA ON CACHE BOOL "") +SET(Kokkos_ARCH_AMPERE80 ON CACHE BOOL "") +#SET(Kokkos_ARCH_ZEN2 ON CACHE BOOL "") # works, and perf same if both AMPERE80 and ZEN2 are on +#SET(Kokkos_ENABLE_CUDA_UVM ON CACHE BOOL "") +SET(Kokkos_ENABLE_EXPLICIT_INSTANTIATION OFF CACHE BOOL "") +#SET(Kokkos_ENABLE_CUDA_ARCH_LINKING OFF CACHE BOOL "") + +#SET(CMAKE_C_COMPILER "mpicc" CACHE STRING "") +#SET(CMAKE_Fortran_COMPILER "mpifort" CACHE STRING "") +#SET(CMAKE_CXX_COMPILER "mpicxx" CACHE STRING "") +SET(CMAKE_C_COMPILER "cc" CACHE STRING "") +SET(CMAKE_Fortran_COMPILER "ftn" CACHE STRING "") +SET(CMAKE_CXX_COMPILER "CC" CACHE STRING "") + +#SET(CMAKE_C_COMPILER "mpicc" CACHE STRING "") +#SET(CMAKE_Fortran_COMPILER "mpifort" CACHE STRING "") +#SET(CMAKE_CXX_COMPILER "${CMAKE_CURRENT_SOURCE_DIR}/../../externals/kokkos/bin/nvcc_wrapper" CACHE STRING "") + +# Note: need to set MPICH_CXX env variable and perhaps NVCC_WRAPPER_DEFAULT_COMPILER + +SET(CXXLIB_SUPPORTED_CACHE FALSE CACHE BOOL "") + +SET(ENABLE_OPENMP OFF CACHE BOOL "") +SET(ENABLE_COLUMN_OPENMP OFF CACHE BOOL "") +SET(ENABLE_HORIZ_OPENMP OFF CACHE BOOL "") + +SET(CMAKE_VERBOSE_MAKEFILE ON CACHE BOOL "") + +#SET(HOMME_TESTING_PROFILE "dev" CACHE STRING "") + +SET(USE_NUM_PROCS 4 CACHE STRING "") + +SET(USE_MPIEXEC "srun" CACHE STRING "") +#SET(CPRNC_DIR /global/cfs/cdirs/e3sm/tools/cprnc CACHE FILEPATH "") diff --git a/components/homme/cmake/machineFiles/spot-aot-AB2.cmake b/components/homme/cmake/machineFiles/spot-aot-AB2.cmake new file mode 100644 index 000000000000..23fad2361ccf --- /dev/null +++ b/components/homme/cmake/machineFiles/spot-aot-AB2.cmake @@ -0,0 +1,63 @@ +#module restore +#module load oneapi/eng-compiler/2022.12.30.005 +#module load intel_compute_runtime/release/agama-devel-627 +#module load spack cmake +#module list + +SET (SUNSPOT_MACHINE TRUE CACHE BOOL "") + +SET (HOMMEXX_MPI_ON_DEVICE TRUE CACHE BOOL "") + +#SET(BUILD_HOMME_WITHOUT_PIOLIBRARY TRUE CACHE BOOL "") + +SET(HOMME_FIND_BLASLAPACK TRUE CACHE BOOL "") + +SET(WITH_PNETCDF FALSE CACHE FILEPATH "") + +SET(USE_QUEUING FALSE CACHE BOOL "") + +#temp hack +SET(HOMME_USE_KOKKOS TRUE CACHE BOOL "") + +SET(BUILD_HOMME_PREQX_KOKKOS TRUE CACHE BOOL "") +SET(BUILD_HOMME_THETA_KOKKOS TRUE CACHE BOOL "") + +#set(KOKKOS_HOME "/home/onguba/kokkos-build/june22-2024-aot/install" CACHE STRING "") +#set(E3SM_KOKKOS_PATH ${KOKKOS_HOME} CACHE STRING "") + +SET (NetCDF_Fortran_PATH "/lus/gila/projects/CSC249ADSE15_CNDA/software/oneAPI.2022.12.30.003/netcdf" CACHE STRING "") +SET (NetCDF_C_PATH "/lus/gila/projects/CSC249ADSE15_CNDA/software/oneAPI.2022.12.30.003/netcdf" CACHE STRING "") + +SET(USE_TRILINOS OFF CACHE BOOL "") + +SET(SYCL_BUILD TRUE CACHE BOOL "") +SET(HOMME_ENABLE_COMPOSE FALSE CACHE BOOL "") + +#SET(CMAKE_CXX_STANDARD 17) +SET(CMAKE_CXX_STANDARD 17 CACHE STRING "CXX Standard") + +SET(CMAKE_C_COMPILER "mpicc" CACHE STRING "") +SET(CMAKE_Fortran_COMPILER "mpifort" CACHE STRING "") +SET(CMAKE_CXX_COMPILER "mpicxx" CACHE STRING "") + +SET(SYCL_COMPILE_FLAGS "-std=c++17 -fsycl -fsycl-device-code-split=per_kernel -fno-sycl-id-queries-fit-in-int -fsycl-unnamed-lambda") +SET(SYCL_LINK_FLAGS "-fsycl-max-parallel-link-jobs=32 -fsycl-link-huge-device-code -fsycl -fsycl-device-code-split=per_kernel -fsycl-targets=spir64_gen -Xsycl-target-backend \"-device 12.60.7\"") + +#-fpscomp does not actually solve the issue with bools in here,another suggestion was -fp-model=precise, not working either +SET(ADD_Fortran_FLAGS " -fc=ifx -fpscomp logicals -O3 -DNDEBUG -DCPRINTEL -g" CACHE STRING "") +SET(ADD_C_FLAGS "-O3 -DNDEBUG " CACHE STRING "") + +SET(ADD_CXX_FLAGS " -std=c++17 -O3 -DNDEBUG ${SYCL_COMPILE_FLAGS}" CACHE STRING "") +SET(ADD_LINKER_FLAGS "-O3 -DNDEBUG ${SYCL_LINK_FLAGS} -fortlib" CACHE STRING "") + +set (ENABLE_OPENMP OFF CACHE BOOL "") +set (ENABLE_COLUMN_OPENMP OFF CACHE BOOL "") +set (ENABLE_HORIZ_OPENMP OFF CACHE BOOL "") + +set (HOMME_TESTING_PROFILE "dev" CACHE STRING "") + +set (USE_NUM_PROCS 4 CACHE STRING "") + +SET (USE_MPI_OPTIONS "--bind-to core" CACHE FILEPATH "") + + diff --git a/components/homme/src/preqx_kokkos/cxx/CamForcing.cpp b/components/homme/src/preqx_kokkos/cxx/CamForcing.cpp index 2b1e6514389e..36ca5f4a95f4 100644 --- a/components/homme/src/preqx_kokkos/cxx/CamForcing.cpp +++ b/components/homme/src/preqx_kokkos/cxx/CamForcing.cpp @@ -51,7 +51,7 @@ void state_forcing( void tracer_forcing( const ExecViewUnmanaged &f_q, const HybridVCoord &hvcoord, const TimeLevel &tl, const int &num_q, - const MoistDry &moisture, const double &dt, + const bool &use_moisture, const double &dt, const ExecViewManaged &ps_v, const ExecViewManaged< Scalar * [Q_NUM_TIME_LEVELS][QSIZE_D][NP][NP][NUM_LEV]> &qdp, @@ -61,7 +61,7 @@ void tracer_forcing( const int np1 = tl.n0; const int np1_qdp = tl.n0_qdp; - if (moisture == MoistDry::MOIST) { + if (use_moisture) { // Remove the m_fq_ps_v buffer since it's not actually needed. // Instead apply the forcing to m_ps_v directly // Bonus - one less parallel reduce in dry cases! @@ -161,7 +161,7 @@ void apply_cam_forcing(const Real &dt) { tracers.fq = decltype(tracers.fq)("fq", elems.num_elems(),tracers.num_tracers()); } tracer_forcing(tracers.fq, hvcoord, tl, tracers.num_tracers(), - sim_params.moisture, dt, elems.m_state.m_ps_v, tracers.qdp, tracers.Q); + sim_params.use_moisture, dt, elems.m_state.m_ps_v, tracers.qdp, tracers.Q); GPTLstop("ApplyCAMForcing"); } diff --git a/components/homme/src/preqx_kokkos/cxx/cxx_f90_interface_preqx.cpp b/components/homme/src/preqx_kokkos/cxx/cxx_f90_interface_preqx.cpp index c75143a9836a..b433a48c2abc 100644 --- a/components/homme/src/preqx_kokkos/cxx/cxx_f90_interface_preqx.cpp +++ b/components/homme/src/preqx_kokkos/cxx/cxx_f90_interface_preqx.cpp @@ -37,7 +37,7 @@ void init_simulation_params_c (const int& remap_alg, const int& limiter_option, const int& time_step_type, const int& qsize, const int& state_frequency, const Real& nu, const Real& nu_p, const Real& nu_q, const Real& nu_s, const Real& nu_div, const Real& nu_top, const int& hypervis_order, const int& hypervis_subcycle, const double& hypervis_scaling, - const int& ftype, const bool& prescribed_wind, const bool& moisture, const bool& disable_diagnostics, + const int& ftype, const bool& prescribed_wind, const bool& use_moisture, const bool& disable_diagnostics, const bool& use_cpstar, const int& transport_alg, const int& dt_remap_factor, const int& dt_tracer_factor, const double& scale_factor, const double& laplacian_rigid_factor) @@ -90,7 +90,7 @@ void init_simulation_params_c (const int& remap_alg, const int& limiter_option, params.hypervis_subcycle = hypervis_subcycle; params.hypervis_scaling = hypervis_scaling; params.disable_diagnostics = disable_diagnostics; - params.moisture = (moisture ? MoistDry::MOIST : MoistDry::DRY); + params.use_moisture = use_moisture; params.use_cpstar = use_cpstar; params.transport_alg = transport_alg; // SphereOperators parameters; preqx supports only the sphere. diff --git a/components/homme/src/preqx_kokkos/cxx/prim_advance_exp.cpp b/components/homme/src/preqx_kokkos/cxx/prim_advance_exp.cpp index f7c7600aab8d..58e58f0160bf 100644 --- a/components/homme/src/preqx_kokkos/cxx/prim_advance_exp.cpp +++ b/components/homme/src/preqx_kokkos/cxx/prim_advance_exp.cpp @@ -34,7 +34,7 @@ void prim_advance_exp (TimeLevel& tl, const Real dt, const bool compute_diagnost // Determine the tracers time level tl.n0_qdp= -1; - if (params.moisture == MoistDry::MOIST) { + if (params.use_moisture) { tl.update_tracers_levels(params.qsplit); } diff --git a/components/homme/src/prim_main.F90 b/components/homme/src/prim_main.F90 index bfbe57e8b317..d6901151d365 100644 --- a/components/homme/src/prim_main.F90 +++ b/components/homme/src/prim_main.F90 @@ -20,7 +20,7 @@ program prim_main use element_mod, only: element_t use common_io_mod, only: output_dir, infilenames use common_movie_mod, only: nextoutputstep - use perf_mod, only: t_initf, t_prf, t_finalizef, t_startf, t_stopf ! _EXTERNAL + use perf_mod, only: t_initf, t_prf, t_finalizef, t_startf, t_stopf, t_disablef, t_enablef ! _EXTERNAL use restart_io_mod , only: restartheader_t, writerestart use hybrid_mod, only: hybrid_create #if (defined MODEL_THETA_L && defined ARKODE) @@ -240,6 +240,11 @@ end subroutine finalize_kokkos_f90 nstep = nextoutputstep(tl) do while(tl%nstep= 2) call t_enablef() call t_startf('prim_run') call prim_run_subcycle(elem, hybrid,nets,nete, tstep, .false., tl, hvcoord,1) call t_stopf('prim_run') diff --git a/components/homme/src/share/cxx/Config.hpp b/components/homme/src/share/cxx/Config.hpp index 684f9143beaf..b204b1dbd047 100644 --- a/components/homme/src/share/cxx/Config.hpp +++ b/components/homme/src/share/cxx/Config.hpp @@ -21,7 +21,7 @@ # endif #endif -#if ! defined HOMMEXX_CUDA_SPACE && ! defined HOMMEXX_OPENMP_SPACE && ! defined HOMMEXX_THREADS_SPACE && ! defined HOMMEXX_SERIAL_SPACE && ! defined HOMMEXX_HIP_SPACE +#if ! defined HOMMEXX_CUDA_SPACE && ! defined HOMMEXX_OPENMP_SPACE && ! defined HOMMEXX_THREADS_SPACE && ! defined HOMMEXX_SERIAL_SPACE && ! defined HOMMEXX_HIP_SPACE && ! defined HOMMEXX_SYCL_SPACE # define HOMMEXX_DEFAULT_SPACE #endif diff --git a/components/homme/src/share/cxx/EulerStepFunctorImpl.hpp b/components/homme/src/share/cxx/EulerStepFunctorImpl.hpp index f3029764dac3..f87bb108bebf 100644 --- a/components/homme/src/share/cxx/EulerStepFunctorImpl.hpp +++ b/components/homme/src/share/cxx/EulerStepFunctorImpl.hpp @@ -652,7 +652,10 @@ class EulerStepFunctorImpl { minmax_and_biharmonic(); } } + + GPTLstart("tl-at adv-n-limit"); advect_and_limit(); + GPTLstop("tl-at adv-n-limit"); exchange_qdp_dss_var(); } @@ -667,6 +670,7 @@ class EulerStepFunctorImpl { void run_tracer_phase (const KernelVariables& kv) const { compute_qtens(kv); kv.team_barrier(); + if (m_data.limiter_option == 8) { limiter_optim_iter_full(kv); kv.team_barrier(); @@ -674,6 +678,7 @@ class EulerStepFunctorImpl { limiter_clip_and_sum(kv); kv.team_barrier(); } + apply_spheremp(kv); } diff --git a/components/homme/src/share/cxx/ExecSpaceDefs.cpp b/components/homme/src/share/cxx/ExecSpaceDefs.cpp index 8d496bff5d16..4f3d97135fea 100644 --- a/components/homme/src/share/cxx/ExecSpaceDefs.cpp +++ b/components/homme/src/share/cxx/ExecSpaceDefs.cpp @@ -21,6 +21,10 @@ #include #endif +#ifdef KOKKOS_ENABLE_SYCL +#include +#endif + namespace Homme { // Since we're initializing from inside a Fortran code and don't have access to @@ -52,7 +56,16 @@ void initialize_kokkos () { // It isn't a big deal if we can't get the device count. nd = 1; } +#elif defined(KOKKOS_ENABLE_SYCL) + +//https://developer.codeplay.com/products/computecpp/ce/2.11.0/guides/sycl-for-cuda-developers/migrating-from-cuda-to-sycl + +//to make it build + int nd = 1; + #endif + + #ifdef HOMMEXX_ENABLE_GPU std::stringstream ss; ss << "--kokkos-num-devices=" << nd; @@ -117,6 +130,7 @@ team_num_threads_vectors_for_gpu ( assert(num_warps_total >= max_num_warps); assert(tp.max_threads_usable >= 1 && tp.max_vectors_usable >= 1); +#ifndef KOKKOS_ENABLE_SYCL int num_warps; if (tp.prefer_larger_team) { const int num_warps_usable = @@ -161,6 +175,9 @@ team_num_threads_vectors_for_gpu ( return std::make_pair( num_device_threads / num_vectors, num_vectors ); } +#else + return std::make_pair(16,8); +#endif } } // namespace Parallel diff --git a/components/homme/src/share/cxx/ExecSpaceDefs.hpp b/components/homme/src/share/cxx/ExecSpaceDefs.hpp index cd6649c7ab2d..82f5e803801c 100644 --- a/components/homme/src/share/cxx/ExecSpaceDefs.hpp +++ b/components/homme/src/share/cxx/ExecSpaceDefs.hpp @@ -34,6 +34,10 @@ using HommexxGPU = Kokkos::Cuda; using HommexxGPU = Kokkos::Experimental::HIP; #endif +#ifdef KOKKOS_ENABLE_SYCL +using HommexxGPU = Kokkos::Experimental::SYCL; +#endif + #else using HommexxGPU = void; #endif diff --git a/components/homme/src/share/cxx/GllFvRemap.cpp b/components/homme/src/share/cxx/GllFvRemap.cpp index e36dbc14d74f..7b0400427f38 100644 --- a/components/homme/src/share/cxx/GllFvRemap.cpp +++ b/components/homme/src/share/cxx/GllFvRemap.cpp @@ -16,13 +16,13 @@ namespace Homme { void init_gllfvremap_c (int nelemd, int np, int nf, int nf_max, - bool theta_hydrostatic_mode, + int theta_hydrostatic_mode, CF90Ptr fv_metdet, CF90Ptr g2f_remapd, CF90Ptr f2g_remapd, CF90Ptr D_f, CF90Ptr Dinv_f) { auto& c = Context::singleton(); auto& g = c.get(); - g.init_data(nf, nf_max, theta_hydrostatic_mode, fv_metdet, g2f_remapd, - f2g_remapd, D_f, Dinv_f); + const bool thm = static_cast(theta_hydrostatic_mode); + g.init_data(nf, nf_max, thm, fv_metdet, g2f_remapd, f2g_remapd, D_f, Dinv_f); } GllFvRemap::GllFvRemap () { @@ -52,7 +52,7 @@ void GllFvRemap::init_boundary_exchanges () { } void GllFvRemap -::init_data (const int nf, const int nf_max, bool theta_hydrostatic_mode, +::init_data (const int nf, const int nf_max, const bool theta_hydrostatic_mode, const Real* fv_metdet, const Real* g2f_remapd, const Real* f2g_remapd, const Real* D_f, const Real* Dinv_f) { m_impl->init_data(nf, nf_max, theta_hydrostatic_mode, fv_metdet, diff --git a/components/homme/src/share/cxx/GllFvRemap.hpp b/components/homme/src/share/cxx/GllFvRemap.hpp index 07e4bf58a903..7ebf5a82b71a 100644 --- a/components/homme/src/share/cxx/GllFvRemap.hpp +++ b/components/homme/src/share/cxx/GllFvRemap.hpp @@ -40,7 +40,7 @@ class GllFvRemap { typedef Phys2T::const_type CPhys2T; typedef Phys3T::const_type CPhys3T; - void init_data(const int nf, const int nf_max, bool theta_hydrostatic_mode, + void init_data(const int nf, const int nf_max, const bool theta_hydrostatic_mode, const Real* fv_metdet, const Real* g2f_remapd, const Real* f2g_remapd, const Real* D_f, const Real* Dinv_f); @@ -81,7 +81,7 @@ class GllFvRemap { extern "C" void init_gllfvremap_c(int nelemd, int np, int nf, int nf_max, - const bool theta_hydrostatic_mode, + const int theta_hydrostatic_mode, CF90Ptr fv_metdet, CF90Ptr g2f_remapd, CF90Ptr f2g_remapd, CF90Ptr D_f, CF90Ptr Dinv_f); diff --git a/components/homme/src/share/cxx/GllFvRemapImpl.cpp b/components/homme/src/share/cxx/GllFvRemapImpl.cpp index 6148f69cfa9c..ea1a52f5efdf 100644 --- a/components/homme/src/share/cxx/GllFvRemapImpl.cpp +++ b/components/homme/src/share/cxx/GllFvRemapImpl.cpp @@ -142,7 +142,7 @@ ::init_data (const int nf, const int nf_max, const bool theta_hydrostatic_mode, " nf must be > 1.", Errors::err_not_implemented); auto& sp = Context::singleton().get(); - m_data.use_moisture = sp.moisture == MoistDry::MOIST; + m_data.use_moisture = sp.use_moisture; // Only in the unit test gllfvremap_ut does theta_hydrostatic_mode not already // == sp.theta_hydrostatic_mode. m_data.theta_hydrostatic_mode = sp.theta_hydrostatic_mode = theta_hydrostatic_mode; diff --git a/components/homme/src/share/cxx/HommexxEnums.hpp b/components/homme/src/share/cxx/HommexxEnums.hpp index 59c8f3c9652c..06abbf35adbc 100644 --- a/components/homme/src/share/cxx/HommexxEnums.hpp +++ b/components/homme/src/share/cxx/HommexxEnums.hpp @@ -47,11 +47,6 @@ enum class ForcingAlg : int { FORCING_2 = 2, // TODO: Rename FORCING_1 and FORCING_2 to something more descriptive }; -enum class MoistDry { - MOIST, - DRY -}; - enum class AdvectionForm { Conservative, NonConservative diff --git a/components/homme/src/share/cxx/SimulationParams.hpp b/components/homme/src/share/cxx/SimulationParams.hpp index b435911da2e6..4f36962b16c3 100644 --- a/components/homme/src/share/cxx/SimulationParams.hpp +++ b/components/homme/src/share/cxx/SimulationParams.hpp @@ -23,7 +23,7 @@ struct SimulationParams void print(std::ostream& out = std::cout); TimeStepType time_step_type; - MoistDry moisture; + bool use_moisture; RemapAlg remap_alg; TestCase test_case; ForcingAlg ftype = ForcingAlg::FORCING_OFF; @@ -77,7 +77,7 @@ inline void SimulationParams::print (std::ostream& out) { out << "\n************** CXX SimulationParams **********************\n\n"; out << " time_step_type: " << etoi(time_step_type) << "\n"; - out << " moisture: " << (moisture==MoistDry::DRY ? "dry" : "moist") << "\n"; + out << " use_moisture: " << (use_moisture ? "moist" : "dry") << "\n"; out << " remap_alg: " << etoi(remap_alg) << "\n"; out << " test case: " << etoi(test_case) << "\n"; out << " ftype: " << etoi(ftype) << "\n"; diff --git a/components/homme/src/share/cxx/utilities/BfbUtils.hpp b/components/homme/src/share/cxx/utilities/BfbUtils.hpp index e3570874e266..7fb4d042f7f2 100644 --- a/components/homme/src/share/cxx/utilities/BfbUtils.hpp +++ b/components/homme/src/share/cxx/utilities/BfbUtils.hpp @@ -64,7 +64,7 @@ KOKKOS_INLINE_FUNCTION ScalarType int_pow (ScalarType val, int k) { constexpr int max_shift = 30; if (k<0) { - printf ("k = %d\n",k); + Kokkos::printf ("k = %d\n",k); Kokkos::abort("int_pow implemented only for k>=0.\n"); } diff --git a/components/homme/src/share/cxx/utilities/scream_tridiag.hpp b/components/homme/src/share/cxx/utilities/scream_tridiag.hpp index e18bbc4e7e27..26221db39552 100644 --- a/components/homme/src/share/cxx/utilities/scream_tridiag.hpp +++ b/components/homme/src/share/cxx/utilities/scream_tridiag.hpp @@ -128,6 +128,10 @@ int get_thread_id_within_team_gpu (const TeamMember& team) { // Can't use team.team_rank() here because vector direction also uses physical // threads but TeamMember types don't expose that information. return blockDim.x * threadIdx.y + threadIdx.x; +#elif defined(__SYCL_DEVICE_ONLY__) + auto item = team.item(); + return static_cast(item.get_local_range(1) * item.get_local_id(0) + + item.get_local_id(1)); #else assert(0); return -1; @@ -138,6 +142,9 @@ template KOKKOS_FORCEINLINE_FUNCTION int get_team_nthr_gpu (const TeamMember& team) { #if defined __CUDA_ARCH__ || defined __HIP_DEVICE_COMPILE__ return blockDim.x * blockDim.y; +#elif defined __SYCL_DEVICE_ONLY__ + auto item = team.item(); + return static_cast(item.get_local_range(0) * item.get_local_range(1)); #else assert(0); return -1; @@ -161,6 +168,16 @@ KOKKOS_FORCEINLINE_FUNCTION int get_team_nthr (const Kokkos::Impl::HIPTeamMember& team) { return get_team_nthr_gpu(team); } #endif // KOKKOS_ENABLE_HIP + +#ifdef KOKKOS_ENABLE_SYCL +KOKKOS_FORCEINLINE_FUNCTION +int get_thread_id_within_team (const Kokkos::Impl::SYCLTeamMember& team) +{ return get_thread_id_within_team_gpu(team); } +KOKKOS_FORCEINLINE_FUNCTION +int get_team_nthr (const Kokkos::Impl::SYCLTeamMember& team) +{ return get_team_nthr_gpu(team); } +#endif // KOKKOS_ENABLE_SYCL + template KOKKOS_INLINE_FUNCTION const T& min (const T& a, const T& b) { return a < b ? a : b; } @@ -634,7 +651,7 @@ void bfb (const TeamMember& team, const auto f = [&] (const int& j) { impl::bfb_thomas_solve(dl, d, du, Kokkos::subview(X , Kokkos::ALL(), j)); }; - Kokkos::parallel_for(Kokkos::TeamThreadRange(team, nrhs), f); + Kokkos::parallel_for(Kokkos::TeamVectorRange(team, nrhs), f); } template @@ -664,7 +681,7 @@ void bfb (const TeamMember& team, subview(du, ALL(), j), subview(X , ALL(), j)); }; - Kokkos::parallel_for(Kokkos::TeamThreadRange(team, nrhs), f); + Kokkos::parallel_for(Kokkos::TeamVectorRange(team, nrhs), f); } } // namespace tridiag diff --git a/components/homme/src/share/gllfvremap_mod.F90 b/components/homme/src/share/gllfvremap_mod.F90 index e8013df217b1..d9dd8e33365a 100644 --- a/components/homme/src/share/gllfvremap_mod.F90 +++ b/components/homme/src/share/gllfvremap_mod.F90 @@ -266,21 +266,22 @@ end subroutine gfr_init subroutine gfr_init_hxx() bind(c) #if KOKKOS_TARGET use control_mod, only: theta_hydrostatic_mode - use iso_c_binding, only: c_bool + use iso_c_binding, only: c_int interface subroutine init_gllfvremap_c(nelemd, np, nf, nf_max, theta_hydrostatic_mode, & fv_metdet, g2f_remapd, f2g_remapd, D_f, Dinv_f) bind(c) - use iso_c_binding, only: c_bool, c_int, c_double + use iso_c_binding, only: c_int, c_double integer (c_int), value, intent(in) :: nelemd, np, nf, nf_max - logical (c_bool), value, intent(in) :: theta_hydrostatic_mode + integer (c_int), value, intent(in) :: theta_hydrostatic_mode real (c_double), dimension(nf*nf,nelemd), intent(in) :: fv_metdet real (c_double), dimension(np,np,nf_max*nf_max), intent(in) :: g2f_remapd real (c_double), dimension(nf_max*nf_max,np,np), intent(in) :: f2g_remapd real (c_double), dimension(nf*nf,2,2,nelemd), intent(in) :: D_f, Dinv_f end subroutine init_gllfvremap_c end interface - logical (c_bool) :: thm - thm = theta_hydrostatic_mode + integer (c_int) :: thm + thm = 0 + if (theta_hydrostatic_mode) thm = 1 call init_gllfvremap_c(nelemd, np, gfr%nphys, nphys_max, thm, & gfr%fv_metdet, gfr%g2f_remapd, gfr%f2g_remapd, gfr%D_f, gfr%Dinv_f) #endif diff --git a/components/homme/src/theta-l_kokkos/cxx/CamForcing.cpp b/components/homme/src/theta-l_kokkos/cxx/CamForcing.cpp index 02b999db16e9..bd7cee3e7c0a 100644 --- a/components/homme/src/theta-l_kokkos/cxx/CamForcing.cpp +++ b/components/homme/src/theta-l_kokkos/cxx/CamForcing.cpp @@ -33,7 +33,7 @@ static void apply_cam_forcing_tracers(const Real dt, ForcingFunctor& ff, if ( p.ftype == ForcingAlg::FORCING_2) adjustment = true; #endif - ff.tracers_forcing(dt, tl.n0, tl.n0_qdp, adjustment, p.moisture); + ff.tracers_forcing(dt, tl.n0, tl.n0_qdp, adjustment, p.use_moisture); GPTLstop("ApplyCAMForcing_tracers"); } diff --git a/components/homme/src/theta-l_kokkos/cxx/DirkFunctorImpl.hpp b/components/homme/src/theta-l_kokkos/cxx/DirkFunctorImpl.hpp index ace1ba920141..d16769079729 100644 --- a/components/homme/src/theta-l_kokkos/cxx/DirkFunctorImpl.hpp +++ b/components/homme/src/theta-l_kokkos/cxx/DirkFunctorImpl.hpp @@ -382,8 +382,8 @@ struct DirkFunctorImpl { kv.team_barrier(); if (it >= maxiter) { - printf("[DIRK] WARNING! Newton reached max iteration count," - " with deltaerr = %3.17f\n", deltaerr); + Kokkos::printf("[DIRK] WARNING! Newton reached max iteration count," + " with deltaerr = %3.17f\n", deltaerr); nerr = 1; } diff --git a/components/homme/src/theta-l_kokkos/cxx/ForcingFunctor.hpp b/components/homme/src/theta-l_kokkos/cxx/ForcingFunctor.hpp index 28a702c1d273..00fa1deef667 100644 --- a/components/homme/src/theta-l_kokkos/cxx/ForcingFunctor.hpp +++ b/components/homme/src/theta-l_kokkos/cxx/ForcingFunctor.hpp @@ -236,7 +236,7 @@ class ForcingFunctor }); } - void tracers_forcing (const Real dt, const int np1, const int np1_qdp, const bool adjustment, const MoistDry moisture) { + void tracers_forcing (const Real dt, const int np1, const int np1_qdp, const bool adjustment, const bool use_moisture) { // The Functor needs to be fully setup to use this function assert (is_setup); @@ -245,7 +245,7 @@ class ForcingFunctor m_np1_qdp = np1_qdp; m_adjustment = adjustment; - m_moist = (moisture==MoistDry::MOIST); + m_moist = use_moisture; Kokkos::parallel_for("temperature, NH perturb press, FQps",m_policy_tracers_pre,*this); Kokkos::fence(); diff --git a/components/homme/src/theta-l_kokkos/cxx/HyperviscosityFunctorImpl.cpp b/components/homme/src/theta-l_kokkos/cxx/HyperviscosityFunctorImpl.cpp index 046e6f9956d4..d160e114475b 100644 --- a/components/homme/src/theta-l_kokkos/cxx/HyperviscosityFunctorImpl.cpp +++ b/components/homme/src/theta-l_kokkos/cxx/HyperviscosityFunctorImpl.cpp @@ -118,9 +118,9 @@ void HyperviscosityFunctorImpl::init_params(const SimulationParams& params) m_eos.init(params.theta_hydrostatic_mode,m_hvcoord); #ifdef HOMMEXX_BFB_TESTING - m_process_nh_vars = true; + m_process_nh_vars = 1; #else - m_process_nh_vars = !params.theta_hydrostatic_mode; + m_process_nh_vars = not params.theta_hydrostatic_mode; #endif } diff --git a/components/homme/src/theta-l_kokkos/cxx/LimiterFunctor.hpp b/components/homme/src/theta-l_kokkos/cxx/LimiterFunctor.hpp index cd3bf7c32526..7914c0a60e3a 100644 --- a/components/homme/src/theta-l_kokkos/cxx/LimiterFunctor.hpp +++ b/components/homme/src/theta-l_kokkos/cxx/LimiterFunctor.hpp @@ -141,8 +141,8 @@ struct LimiterFunctor { [&](const int k,Real& result) { #ifndef HOMMEXX_BFB_TESTING if(diff_as_real(k) < 0){ - printf("WARNING:CAAR: dp3d too small. k=%d, dp3d(k)=%f, dp0=%f \n", - k+1,dp_as_real(k),dp0_as_real(k)); + Kokkos::printf("WARNING:CAAR: dp3d too small. k=%d, dp3d(k)=%f, dp0=%f \n", + k+1,dp_as_real(k),dp0_as_real(k)); } #endif result = result<=diff_as_real(k) ? result : diff_as_real(k); @@ -202,8 +202,8 @@ struct LimiterFunctor { for (int ivec=0; ivec(e.m_forcing); } -void init_functors_c (const bool& allocate_buffer) +void init_functors_c (const int& allocate_buffer) { auto& c = Context::singleton(); diff --git a/components/homme/src/theta-l_kokkos/prim_driver_mod.F90 b/components/homme/src/theta-l_kokkos/prim_driver_mod.F90 index 96b42314453f..eae8544ca865 100644 --- a/components/homme/src/theta-l_kokkos/prim_driver_mod.F90 +++ b/components/homme/src/theta-l_kokkos/prim_driver_mod.F90 @@ -64,7 +64,7 @@ subroutine prim_init2(elem, hybrid, nets, nete, tl, hvcoord) end subroutine prim_init2 subroutine prim_create_c_data_structures (tl, hvcoord, mp) - use iso_c_binding, only : c_loc, c_ptr, c_bool, C_NULL_CHAR + use iso_c_binding, only : c_loc, c_ptr, C_NULL_CHAR use theta_f2c_mod, only : init_reference_element_c, init_simulation_params_c, & init_time_level_c, init_hvcoord_c, init_elements_c use time_mod, only : TimeLevel_t, nsplit @@ -73,7 +73,7 @@ subroutine prim_create_c_data_structures (tl, hvcoord, mp) nu, nu_p, nu_q, nu_s, nu_div, nu_top, vert_remap_q_alg, & hypervis_order, hypervis_subcycle, hypervis_subcycle_tom,& hypervis_scaling, & - ftype, prescribed_wind, moisture, disable_diagnostics, & + ftype, prescribed_wind, use_moisture, disable_diagnostics, & use_cpstar, transport_alg, theta_hydrostatic_mode, & dcmip16_mu, theta_advect_form, test_case, & MAX_STRING_LEN, dt_remap_factor, dt_tracer_factor, & @@ -93,6 +93,8 @@ subroutine prim_create_c_data_structures (tl, hvcoord, mp) type (c_ptr) :: hybrid_am_ptr, hybrid_ai_ptr, hybrid_bm_ptr, hybrid_bi_ptr character(len=MAX_STRING_LEN), target :: test_name + integer :: disable_diagnostics_int, theta_hydrostatic_mode_int, use_moisture_int + ! Initialize the C++ reference element structure (i.e., pseudo-spectral deriv matrix and ref element mass matrix) dvv = deriv1%dvv elem_mp = mp @@ -100,22 +102,30 @@ subroutine prim_create_c_data_structures (tl, hvcoord, mp) ! Fill the simulation params structures in C++ test_name = TRIM(test_case) // C_NULL_CHAR + + disable_diagnostics_int = 0 + if (disable_diagnostics) disable_diagnostics_int = 1 + use_moisture_int = 0 + if (use_moisture) use_moisture_int = 1 + theta_hydrostatic_mode_int = 0 + if (theta_hydrostatic_mode) theta_hydrostatic_mode_int = 1 + call init_simulation_params_c (vert_remap_q_alg, limiter_option, rsplit, qsplit, tstep_type, & qsize, statefreq, nu, nu_p, nu_q, nu_s, nu_div, nu_top, & hypervis_order, hypervis_subcycle, hypervis_subcycle_tom, & hypervis_scaling, & dcmip16_mu, ftype, theta_advect_form, & - LOGICAL(prescribed_wind==1,c_bool), & - LOGICAL(moisture/="dry",c_bool), & - LOGICAL(disable_diagnostics,c_bool), & - LOGICAL(use_cpstar==1,c_bool), & + prescribed_wind, & + use_moisture_int, & + disable_diagnostics_int, & + use_cpstar, & transport_alg, & - LOGICAL(theta_hydrostatic_mode,c_bool), & + theta_hydrostatic_mode_int, & c_loc(test_name), & dt_remap_factor, dt_tracer_factor, & scale_factor, laplacian_rigid_factor, & nsplit, & - LOGICAL(pgrad_correction==1,c_bool), & + pgrad_correction, & dp3d_thresh, vtheta_thresh, internal_diagnostics_level) ! Initialize time level structure in C++ @@ -343,22 +353,21 @@ subroutine prim_init_elements_views (elem) end subroutine prim_init_elements_views subroutine prim_init_kokkos_functors (allocate_buffer) - use iso_c_binding, only : c_bool + use iso_c_binding, only : c_int use theta_f2c_mod, only : init_functors_c, init_boundary_exchanges_c - ! ! Optional Input ! - logical(kind=c_bool), optional :: allocate_buffer ! Whether functor memory buffer should be allocated internally - + logical, intent(in), optional :: allocate_buffer ! Whether functor memory buffer should be allocated internally + integer(kind=c_int) :: ab ! Initialize the C++ functors in the C++ context ! If no argument allocate_buffer is present, ! let Homme internally allocate buffers + ab = 1 if (present(allocate_buffer)) then - call init_functors_c (logical(allocate_buffer,c_bool)) - else - call init_functors_c (logical(.true.,c_bool)) - endif + if (.not. allocate_buffer) ab = 0 + end if + call init_functors_c (ab) ! Initialize boundary exchange structure in C++ call init_boundary_exchanges_c () diff --git a/components/homme/src/theta-l_kokkos/theta_f2c_mod.F90 b/components/homme/src/theta-l_kokkos/theta_f2c_mod.F90 index 7a4c0424807b..ba39bb03c22b 100644 --- a/components/homme/src/theta-l_kokkos/theta_f2c_mod.F90 +++ b/components/homme/src/theta-l_kokkos/theta_f2c_mod.F90 @@ -11,14 +11,14 @@ subroutine init_simulation_params_c (remap_alg, limiter_option, rsplit, qsplit, qsize, state_frequency, nu, nu_p, nu_q, nu_s, nu_div, nu_top, & hypervis_order, hypervis_subcycle, hypervis_subcycle_tom, & hypervis_scaling, & - dcmip16_mu, ftype, theta_adv_form, prescribed_wind, moisture, & + dcmip16_mu, ftype, theta_adv_form, prescribed_wind, use_moisture, & disable_diagnostics, use_cpstar, transport_alg, & theta_hydrostatic_mode, test_case_name, dt_remap_factor, & dt_tracer_factor, scale_factor, laplacian_rigid_factor, & nsplit, pgrad_correction, dp3d_thresh, vtheta_thresh, & internal_diagnostics_level) bind(c) - use iso_c_binding, only: c_int, c_bool, c_double, c_ptr + use iso_c_binding, only: c_int, c_double, c_ptr ! ! Inputs ! @@ -29,8 +29,8 @@ subroutine init_simulation_params_c (remap_alg, limiter_option, rsplit, qsplit, scale_factor, laplacian_rigid_factor, dp3d_thresh, vtheta_thresh integer(kind=c_int), intent(in) :: hypervis_order, hypervis_subcycle, hypervis_subcycle_tom integer(kind=c_int), intent(in) :: ftype, theta_adv_form - logical(kind=c_bool), intent(in) :: prescribed_wind, moisture, disable_diagnostics, use_cpstar - logical(kind=c_bool), intent(in) :: theta_hydrostatic_mode, pgrad_correction + integer(kind=c_int), intent(in) :: prescribed_wind, use_moisture, disable_diagnostics, use_cpstar + integer(kind=c_int), intent(in) :: theta_hydrostatic_mode, pgrad_correction type(c_ptr), intent(in) :: test_case_name end subroutine init_simulation_params_c @@ -138,11 +138,11 @@ end subroutine init_reference_element_c ! Create C++ functors subroutine init_functors_c (allocate_buffer) bind(c) - use iso_c_binding, only: c_bool + use iso_c_binding, only: c_int ! ! Inputs ! - logical(kind=c_bool), intent(in) :: allocate_buffer + integer(kind=c_int), intent(in) :: allocate_buffer end subroutine init_functors_c ! Initialize C++ boundary exchange structures diff --git a/components/homme/test_execs/CMakeLists.txt b/components/homme/test_execs/CMakeLists.txt index a3113921b029..a007a5532b61 100644 --- a/components/homme/test_execs/CMakeLists.txt +++ b/components/homme/test_execs/CMakeLists.txt @@ -142,8 +142,11 @@ ADD_CUSTOM_TARGET(test-execs) ADD_CUSTOM_TARGET(check COMMAND ${CMAKE_CTEST_COMMAND} "--output-on-failure") +if(NOT BUILD_HOMME_WITHOUT_PIOLIBRARY) # Force cprnc to be built when make check is run ADD_DEPENDENCIES(check cprnc) +endif() + # Create a target for making the reference data ADD_CUSTOM_TARGET(baseline diff --git a/components/homme/test_execs/share_kokkos_ut/CMakeLists.txt b/components/homme/test_execs/share_kokkos_ut/CMakeLists.txt index 3fbeff9f6f2c..bc788462ce6e 100644 --- a/components/homme/test_execs/share_kokkos_ut/CMakeLists.txt +++ b/components/homme/test_execs/share_kokkos_ut/CMakeLists.txt @@ -10,7 +10,7 @@ SET(UTILS_TIMING_DIRS ${UTILS_TIMING_SRC_DIR} ${UTILS_TIMING_BIN_DIR}) # Note: need CUDA_BUILD and HOMMEXX_BFB_TESTING here, since the share # unit tests do not include a config.h file SET (COMMON_DEFINITIONS NP=4 NC=4) -IF (CUDA_BUILD OR HIP_BUILD) +IF (CUDA_BUILD OR HIP_BUILD OR SYCL_BUILD) SET(COMMON_DEFINITIONS ${COMMON_DEFINITIONS} HOMMEXX_ENABLE_GPU_F90) ENDIF() IF (HOMMEXX_BFB_TESTING) @@ -158,7 +158,7 @@ ELSE() SET (NUM_CPUS 1) ENDIF() cxx_unit_test (sphere_op_ut "${SPHERE_OP_UT_F90_SRCS}" "${SPHERE_OP_UT_CXX_SRCS}" "${SPHERE_OP_UT_INCLUDE_DIRS}" "${CONFIG_DEFINES}" ${NUM_CPUS}) -endif () +endif () #BFB ### Limiters unit test ### diff --git a/components/homme/test_execs/thetal_kokkos_ut/CMakeLists.txt b/components/homme/test_execs/thetal_kokkos_ut/CMakeLists.txt index 205635e918cc..e8bf5e20bd03 100644 --- a/components/homme/test_execs/thetal_kokkos_ut/CMakeLists.txt +++ b/components/homme/test_execs/thetal_kokkos_ut/CMakeLists.txt @@ -11,6 +11,8 @@ SET(UTILS_TIMING_BIN_DIR ${HOMME_BINARY_DIR}/utils/cime/CIME/non_py/src/timing) THETAL_KOKKOS_SETUP() # This is needed to compile the lib and test executables with the correct options +#these vars shared between all targets, so changing one var +#for one test only won't work, config is built once and for the last test SET(THIS_CONFIG_IN ${HOMME_SOURCE_DIR}/src/theta-l_kokkos/config.h.cmake.in) SET(THIS_CONFIG_HC ${CMAKE_CURRENT_BINARY_DIR}/config.h.c) SET(THIS_CONFIG_H ${CMAKE_CURRENT_BINARY_DIR}/config.h) @@ -18,6 +20,7 @@ SET (NUM_POINTS 4) SET (NUM_PLEV 12) SET (QSIZE_D 4) SET (PIO_INTERP TRUE) + HommeConfigFile (${THIS_CONFIG_IN} ${THIS_CONFIG_HC} ${THIS_CONFIG_H} ) ADD_LIBRARY(thetal_kokkos_ut_lib diff --git a/components/homme/test_execs/thetal_kokkos_ut/forcing_ut.cpp b/components/homme/test_execs/thetal_kokkos_ut/forcing_ut.cpp index 5e4c51c7ca11..fb301166f429 100644 --- a/components/homme/test_execs/thetal_kokkos_ut/forcing_ut.cpp +++ b/components/homme/test_execs/thetal_kokkos_ut/forcing_ut.cpp @@ -160,8 +160,8 @@ TEST_CASE("forcing", "forcing") { std::cout << "Testing tracers forcing.\n"; for (const bool hydrostatic : {true,false}) { std::cout << " -> hydrostatic mode: " << (hydrostatic ? "true" : "false") << "\n"; - for (const MoistDry moisture : {MoistDry::DRY,MoistDry::MOIST}) { - std::cout << " -> moisture: " << (moisture==MoistDry::MOIST ? "moist" : "dry") << "\n"; + for (const bool use_moisture: {false,true}) { + std::cout << " -> moisture: " << (use_moisture ? "moist" : "dry") << "\n"; for (const bool adjustment : {true,false}) { std::cout << " -> adjustment: " << (adjustment ? "true" : "false") << "\n"; @@ -200,8 +200,8 @@ TEST_CASE("forcing", "forcing") { ff.init_buffers(fbm); // Run tracers forcing (cxx and f90) - ff.tracers_forcing(dt,np1,np1_qdp,adjustment,moisture); - tracers_forcing_f90(dt,np1+1,np1_qdp+1,hydrostatic,moisture==MoistDry::MOIST,adjustment); + ff.tracers_forcing(dt,np1,np1_qdp,adjustment,use_moisture); + tracers_forcing_f90(dt,np1+1,np1_qdp+1,hydrostatic,use_moisture,adjustment); // Compare answers Kokkos::deep_copy(h_dp,state.m_dp3d); diff --git a/components/homme/test_execs/thetal_kokkos_ut/gllfvremap_ut.cpp b/components/homme/test_execs/thetal_kokkos_ut/gllfvremap_ut.cpp index 0f14b0c3e55a..cf9db941ea1c 100644 --- a/components/homme/test_execs/thetal_kokkos_ut/gllfvremap_ut.cpp +++ b/components/homme/test_execs/thetal_kokkos_ut/gllfvremap_ut.cpp @@ -183,7 +183,7 @@ struct Session { p.qsize = qsize; p.hypervis_scaling = 0; p.transport_alg = 0; - p.moisture = MoistDry::MOIST; + p.use_moisture = true; p.theta_hydrostatic_mode = false; p.scale_factor = is_sphere ? PhysicalConstants::rearth0 : 1; p.laplacian_rigid_factor = is_sphere ? 1/p.scale_factor : 0; @@ -725,7 +725,7 @@ static void test_get_temperature (Session& s) { const auto& sp = c.get(); EquationOfState eos; eos.init(theta_hydrostatic_mode, s.h); ElementOps ops; ops.init(s.h); - const bool use_moisture = sp.moisture == MoistDry::MOIST; + const bool use_moisture = sp.use_moisture; const auto state = c.get(); const auto tracers = c.get(); const auto dp3d = state.m_dp3d;