improving cmake usage for HIP version, minor changes and refactoring

CMakeLists.txt

@@ -56,6 +56,20 @@ if( HAVE_CUDA )
     enable_language( CUDA )
 endif()
 
+ecbuild_add_option( FEATURE HIP
+	DESCRIPTION "HIP" DEFAULT OFF
+	REQUIRED_PACKAGES "hip" )
+if ( HAVE_HIP )
+  if(NOT DEFINED ROCM_PATH)
+    if(DEFINED ENV{ROCM_PATH})
+      set(ROCM_PATH $ENV{ROCM_PATH} CACHE PATH "Path to which ROCM has been installed")
+    else()
+      set(ROCM_PATH "/opt/rocm" CACHE PATH "Path to which ROCM has been installed")
+    endif()
+  endif()
+  find_package(hip REQUIRED)
+endif()
+
 ### OpenMP
 ecbuild_add_option( FEATURE OMP
     DESCRIPTION "OpenMP" DEFAULT ON
@@ -83,7 +97,7 @@ endif()
 ecbuild_add_option( FEATURE SERIALBOX
     DESCRIPTION "Use Serialbox to read input and reference data"
     REQUIRED_PACKAGES "Serialbox"
-    CONDITION NOT HAVE_HDF5 OR HAVE_SERIALBOX
+    CONDITION NOT HAVE_HDF5
     DEFAULT OFF )
 if( HAVE_SERIALBOX )
     list(APPEND CLOUDSC_DEFINITIONS HAVE_SERIALBOX)

arch/eurohpc/lumi/cray-gpu/15.0.1/toolchain.cmake

@@ -18,8 +18,14 @@ set( ENABLE_USE_STMT_FUNC ON CACHE STRING "" )
 ####################################################################
 
 set( ENABLE_OMP ON CACHE STRING "" )
-set( OpenMP_C_FLAGS   "-homp" CACHE STRING "" )
-set( OpenMP_Fortran_FLAGS   "-homp -hnoacc -hlist=aimd -maxrregcount 64" CACHE STRING "" )
+set( OpenMP_C_FLAGS   "-fopenmp" CACHE STRING "" )
+set( OpenMP_CXX_FLAGS "-fopenmp" CACHE STRING "" )
+set( OpenMP_Fortran_FLAGS   "-fopenmp -hnoacc -hlist=aimd" CACHE STRING "" )
+
+set( OpenMP_C_LIB_NAMES       "craymp" )
+set( OpenMP_CXX_LIB_NAMES     "craymp" )
+set( OpenMP_Fortran_LIB_NAMES "craymp" )
+set( OpenMP_craymp_LIBRARY    "/opt/cray/pe/cce/15.0.1/cce/x86_64/lib/libcraymp.so" )
 
 ####################################################################
 # OpenACC FLAGS
@@ -30,6 +36,15 @@ set( OpenACC_C_FLAGS "-hacc" )
 set( OpenACC_CXX_FLAGS "-hacc" )
 set( OpenACC_Fortran_FLAGS "-hacc -h acc_model=deep_copy" )
 
+####################################################################
+# OpenACC FLAGS
+####################################################################
+
+set(CMAKE_HIP_FLAGS "${CMAKE_HIP_FLAGS} -03 -ffast-math")
+if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
+  set(CMAKE_HIP_ARCHITECTURES gfx90a)
+endif()
+
 ####################################################################
 # Compiler FLAGS
 ####################################################################

bundle.yml

@@ -78,10 +78,11 @@ options :
             ENABLE_CLOUDSC_GPU_SCC_CUF=ON
             ENABLE_CLOUDSC_GPU_SCC_CUF_K_CACHING=ON
             BUILD_field_api=ON
+    
     - with-hip :
         help: Enable GPU kernel variant based on HIP
         cmake: >
-            ENABLE_CLOUDSC_HIP=ON 
+            ENABLE_HIP=ON 
 
     - with-mpi :
         help  : Enable MPI-parallel kernel

src/cloudsc_hip/CMakeLists.txt

@@ -8,26 +8,14 @@
 
 # Define this dwarf variant as an ECBuild feature
 ecbuild_add_option( FEATURE CLOUDSC_HIP
-	DESCRIPTION "Build the HIP version CLOUDSC using Serialbox" DEFAULT OFF
-        CONDITION Serialbox_FOUND
+	DESCRIPTION "Build the HIP version CLOUDSC using Serialbox" DEFAULT ON
+	CONDITION Serialbox_FOUND AND HAVE_HIP
 )
 
 if( HAVE_CLOUDSC_HIP )
 
-
-    if(NOT DEFINED ROCM_PATH)
-      if(DEFINED ENV{ROCM_PATH})
-        set(ROCM_PATH $ENV{ROCM_PATH} CACHE PATH "Path to which ROCM has been installed")
-      else()
-        set(ROCM_PATH "/opt/rocm" CACHE PATH "Path to which ROCM has been installed")
-      endif()
-    endif()
-
-    find_package(hip REQUIRED)
-
     set(CMAKE_C_COMPILER "${ROCM_PATH}/bin/hipcc")
     set(CMAKE_CXX_COMPILER "${ROCM_PATH}/bin/hipcc")
-    # set(CMAKE_HIP_FLAGS "${CMAKE_HIP_FLAGS} -03 -ffast-math")
 
     ###### SCC-HIP ####
     ecbuild_add_library(
@@ -57,7 +45,11 @@ if( HAVE_CLOUDSC_HIP )
     target_include_directories(dwarf-cloudsc-hip-lib PUBLIC $<INSTALL_INTERFACE:include> $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/cloudsc>)
     target_link_libraries(dwarf-cloudsc-hip-lib PUBLIC hip::device Serialbox::Serialbox_C $<${HAVE_OMP}:OpenMP::OpenMP_C>)
 
-    target_compile_options(dwarf-cloudsc-hip-lib PRIVATE --offload-arch=gfx90a)
+    if (NOT DEFINED CMAKE_HIP_ARCHITECTURES)
+      message(WARNING "No HIP architecture is set! ('CMAKE_HIP_ARCHITECTURES' is not defined)")
+    else()
+      target_compile_options(dwarf-cloudsc-hip-lib PRIVATE --offload-arch=${CMAKE_HIP_ARCHITECTURES})
+    endif()
 
     ecbuild_add_executable(
         TARGET  dwarf-cloudsc-hip
@@ -95,8 +87,12 @@ if( HAVE_CLOUDSC_HIP )
     target_include_directories(dwarf-cloudsc-hip-hoist-lib PUBLIC $<INSTALL_INTERFACE:include> $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/cloudsc>)
     target_link_libraries(dwarf-cloudsc-hip-hoist-lib PUBLIC hip::device Serialbox::Serialbox_C $<${HAVE_OMP}:OpenMP::OpenMP_C>)
 
-    target_compile_options(dwarf-cloudsc-hip-hoist-lib PRIVATE --offload-arch=gfx90a)
-
+    if (NOT DEFINED CMAKE_HIP_ARCHITECTURES)
+      message(WARNING "No HIP architecture is set! ('CMAKE_HIP_ARCHITECTURES' is not defined)")
+    else()
+      target_compile_options(dwarf-cloudsc-hip-hoist-lib PRIVATE --offload-arch=${CMAKE_HIP_ARCHITECTURES})
+    endif()
+
     ecbuild_add_executable(
         TARGET  dwarf-cloudsc-hip-hoist
         SOURCES dwarf_cloudsc.cpp
@@ -133,8 +129,12 @@ if( HAVE_CLOUDSC_HIP )
     target_include_directories(dwarf-cloudsc-hip-k-caching-lib PUBLIC $<INSTALL_INTERFACE:include> $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/cloudsc>)
     target_link_libraries(dwarf-cloudsc-hip-k-caching-lib PUBLIC hip::device Serialbox::Serialbox_C $<${HAVE_OMP}:OpenMP::OpenMP_C>)
 
-    target_compile_options(dwarf-cloudsc-hip-k-caching-lib PRIVATE --offload-arch=gfx90a)
-
+    if (NOT DEFINED CMAKE_HIP_ARCHITECTURES)
+      message(WARNING "No HIP architecture is set! ('CMAKE_HIP_ARCHITECTURES' is not defined)")
+    else()
+      target_compile_options(dwarf-cloudsc-hip-k-caching-lib PRIVATE --offload-arch=${CMAKE_HIP_ARCHITECTURES})
+    endif()
+
     ecbuild_add_executable(
         TARGET  dwarf-cloudsc-hip-k-caching
         SOURCES dwarf_cloudsc.cpp

src/cloudsc_hip/cloudsc/load_state.cpp

@@ -23,14 +23,13 @@ void query_state(int *klon, int *klev)
 
 void expand_1d(double *buffer, double *field_in, int nlon, int nproma, int ngptot, int nblocks)
 {
-  int b, l, i, buf_start_idx, buf_idx;
+  int b, i, buf_start_idx, buf_idx;
 
-#pragma omp parallel for default(shared) private(b, l, i, buf_start_idx, buf_idx)
+#pragma omp parallel for default(shared) private(b, i, buf_start_idx, buf_idx)
   for (b = 0; b < nblocks; b++) {
     buf_start_idx = ((b)*nproma) % nlon;
     for (i = 0; i < nproma; i++) {
       buf_idx = (buf_start_idx + i) % nlon;
-      // field[b][i] = buffer[buf_idx];
       field_in[b*nproma+i] = buffer[buf_idx];
     }
   }
@@ -39,17 +38,16 @@ void expand_1d(double *buffer, double *field_in, int nlon, int nproma, int ngpto
 
 void expand_1d_int(int *buffer, int *field_in, int nlon, int nproma, int ngptot, int nblocks)
 {
-  int b, l, i, buf_start_idx, buf_idx;
+  int b, i, buf_start_idx, buf_idx;
 
-  #pragma omp parallel for default(shared) private(b, l, i, buf_start_idx, buf_idx)
+  #pragma omp parallel for default(shared) private(b, i, buf_start_idx, buf_idx)
   for (b = 0; b < nblocks; b++) {
     buf_start_idx = ((b)*nproma) % nlon;
     for (i = 0; i < nproma; i++) {
       buf_idx = (buf_start_idx + i) % nlon;
       field_in[b*nproma+i] = buffer[buf_idx];
     }
   }
-
 }
 
 
@@ -61,13 +59,12 @@ void expand_2d(double *buffer_in, double *field_in, int nlon, int nlev, int npro
   for (b = 0; b < nblocks; b++) {
     buf_start_idx = ((b)*nproma) % nlon;
     for (i = 0; i < nproma; i++) {
-    for (l = 0; l < nlev; l++) {
-       buf_idx = (buf_start_idx + i) % nlon;
-       field_in[b*nlev*nproma+l*nproma+i] = buffer_in[l*nlon+buf_idx];
-    }
+      for (l = 0; l < nlev; l++) {
+        buf_idx = (buf_start_idx + i) % nlon;
+        field_in[b*nlev*nproma+l*nproma+i] = buffer_in[l*nlon+buf_idx];
+      }
     }
   }
-
 }
 
 void expand_3d(double *buffer_in, double *field_in, int nlon, int nlev, int nclv, int nproma, int ngptot, int nblocks)
@@ -78,15 +75,14 @@ void expand_3d(double *buffer_in, double *field_in, int nlon, int nlev, int nclv
   for (b = 0; b < nblocks; b++) {
     buf_start_idx = ((b)*nproma) % nlon;
     for (i = 0; i < nproma; i++) {
-    for (c = 0; c < nclv; c++) {
-    for (l = 0; l < nlev; l++) {
-      buf_idx = (buf_start_idx + i) % nlon;
-      field_in[b*nclv*nlev*nproma+c*nlev*nproma+l*nproma+i] = buffer_in[c*nlev*nlon+l*nlon+buf_idx];
-    }
-    }
+      for (c = 0; c < nclv; c++) {
+        for (l = 0; l < nlev; l++) {
+          buf_idx = (buf_start_idx + i) % nlon;
+          field_in[b*nclv*nlev*nproma+c*nlev*nproma+l*nproma+i] = buffer_in[c*nlev*nlon+l*nlon+buf_idx];
+        }
+      }
     }
   }
-
 }