cmake conditions similar to main branch

FastCaloSimAnalyzer/CMakeLists.txt

@@ -23,7 +23,6 @@ set(USE_ALPAKA OFF CACHE BOOL "Use alpaka")
 set(USE_HIP OFF CACHE BOOL "Use HIP")
 set(HIP_TARGET "AMD" CACHE STRING "HIP backend. must be either AMD or NVIDIA")
 
-
 if ( USE_STDPAR )
   if ( ${STDPAR_TARGET} STREQUAL "cpu" )
     if ( NOT RNDGEN_CPU )
@@ -46,11 +45,13 @@ elseif( USE_KOKKOS )
 elseif(USE_ALPAKA)
   find_package(alpaka REQUIRED)
 elseif(USE_HIP)
-  find_package(HIP REQUIRED) 
   if ( ${HIP_TARGET} STREQUAL "NVIDIA" )
+    find_package(HIP) 
     if ( NOT RNDGEN_CPU )
       message(FATAL_ERROR "when HIP_TARGET=NVIDIA, RNDGEN_CPU must be ON")
     endif()
+  else()
+    find_package(HIP REQUIRED) 
   endif()
 endif()
 

FastCaloSimAnalyzer/FastCaloGpu/FastCaloGpu/Rand4Hits.h

@@ -170,6 +170,7 @@ class Rand4Hits {
   unsigned int m_current_hits;
   void *m_gen{ nullptr };
   bool m_useCPU{ false };
+  unsigned long long m_seed{0};
 
   // patch in some GPU pointers for cudaMalloc
   CELL_ENE_T *m_cells_energy{ 0 };

FastCaloSimAnalyzer/FastCaloGpu/src/CMakeLists.txt

@@ -20,37 +20,52 @@ endif()
 if(USE_HIP)
   set(FIND_CUDA OFF)
 endif()
-  
+
 
+if(ENABLE_OMPGPU)
+  string(STRIP ${CMAKE_CXX_FLAGS} OMP_OFFLOAD_TARGET)
+  string(FIND ${CMAKE_CXX_FLAGS} "gfx" OMP_OFFLOAD_TARGET_AMD)
+  string(FIND ${CMAKE_CXX_FLAGS} "sm_" OMP_OFFLOAD_TARGET_NVIDIA)
+  if(OMP_OFFLOAD_TARGET_NVIDIA GREATER 0)
+    message(STATUS "OMP_OFFLOAD_TARGET NVIDIA" )
+  elseif(OMP_OFFLOAD_TARGET_AMD GREATER 0)
+    set(FIND_CUDA OFF)
+    message(STATUS "OMP_OFFLOAD_TARGET AMD" )
+  else()
+    if($ENV{OMP_TARGET_OFFLOAD} MATCHES "disabled")
+      set(FIND_CUDA OFF)
+    else()
+      message(FATAL_ERROR "!! Please specify OpenMP offload target via -DCMAKE_CXX_FLAGS=\"--offload-arch=gfx<>|sm_<>\" or set environment var OMP_TARGET_OFFLOAD=disabled")
+    endif()
+  endif()
+endif()
+
 if(FIND_CUDA)
   find_package(CUDA REQUIRED)
   enable_language( CUDA )
   set(CUDA_LIBRARIES PUBLIC ${CUDA_LIBRARIES})
 endif()
 
-# Add OpenMP
-if(ENABLE_OMPGPU)
-  find_package(OpenMP)
-  if(OPENMP_FOUND)
-  	set(OpenMP_OPT_FLAGS "${OpenMP_OPT_FLAGS} -fopenmp-cuda-mode")
-  	set(OpenMP_OPT_FLAGS "${OpenMP_OPT_FLAGS} -foffload-lto")
-  	set(OpenMP_OPT_FLAGS "${OpenMP_OPT_FLAGS} -fopenmp-assume-no-thread-state")
-  	set(OpenMP_OPT_RMRKS "-Rpass=openmp-opt -Rpass-analysis=openmp-opt -Rpass-missed=openmp-opt " )
-  	set(OpenMP_FLAGS "-fopenmp --offload-arch=sm_86 -lomp") ## nvidia
-  	set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${OpenMP_FLAGS} ${OpenMP_OPT_FLAGS} ${OpenMP_OPT_RMRKS}")
-  	set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OpenMP_FLAGS} ${OpenMP_OPT_FLAGS} ${OpenMP_OPT_RMRKS}")
-  else()
-    message(WARNING "Configuring with OpenMP GPU but OpenMP is not found!")
-  endif()
-endif()
-
 # Sources
 
 if(USE_STDPAR)
   set(FastCaloGpu_Srcs GeoLoadGpu.cxx KernelWrapper_sp.cxx gpuQ.cu Rand4Hits.cu Rand4Hits_sp.cxx )
 elseif(USE_KOKKOS)
   set(FastCaloGpu_Srcs GeoLoadGpu.cxx KernelWrapper_kk.cxx DEV_BigMem_kk.cxx)
 elseif(ENABLE_OMPGPU)
+  # Add OpenMP
+  find_package(OpenMP)
+  if(OPENMP_FOUND)
+    set(OpenMP_OPT_FLAGS "${OpenMP_OPT_FLAGS} -fopenmp-cuda-mode")
+    set(OpenMP_OPT_FLAGS "${OpenMP_OPT_FLAGS} -foffload-lto")
+    set(OpenMP_OPT_FLAGS "${OpenMP_OPT_FLAGS} -fopenmp-assume-no-thread-state")
+    set(OpenMP_OPT_RMRKS "-Rpass=openmp-opt -Rpass-analysis=openmp-opt -Rpass-missed=openmp-opt " )
+    set(OpenMP_FLAGS "-fopenmp -lomp -lomptarget")
+    set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${OpenMP_FLAGS} ${OpenMP_OPT_FLAGS} ${OpenMP_OPT_RMRKS}")
+    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OpenMP_FLAGS} ${OpenMP_OPT_FLAGS} ${OpenMP_OPT_RMRKS}")
+  else()
+    message(WARNING "Configuring with OpenMP GPU but OpenMP is not found!")
+  endif()
   set(FastCaloGpu_Srcs KernelWrapper_omp.cxx gpuQ.cxx CaloGpuGeneral.cxx DEV_BigMem_omp.cxx )
 elseif(USE_ALPAKA)
   set(FastCaloGpu_Srcs CaloGpuGeneral.cxx KernelWrapper_al.cxx Rand4Hits_al.cxx )
@@ -66,10 +81,10 @@ elseif(USE_HIP)
   include_directories( ${ROCM_PATH}/hip/include )
 
   if ( ${HIP_TARGET} STREQUAL "AMD" )
-    message(STATUS "    Using AMD HIP backend")
+    message(STATUS "Using AMD HIP backend")
     set(CMAKE_HIP_ARCHITECTURES "gfx90a;gfx906;gfx908")
   elseif( ${HIP_TARGET} STREQUAL "NVIDIA" )
-    message(STATUS "    Using NVIDIA HIP backend")
+    message(STATUS "Using NVIDIA HIP backend")
     find_package(CUDAToolkit REQUIRED)
     set(CMAKE_CUDA_ARCHITECTURES "70;75;80;86")
     set(CMAKE_HIP_ARCHITECTURES "gfx90a;gfx906;gfx908")
@@ -120,7 +135,20 @@ elseif(USE_HIP)
     target_compile_definitions(${FastCaloGpu_LIB} PUBLIC -DHIP_TARGET_NVIDIA)
     target_link_libraries(${FastCaloGpu_LIB} PUBLIC CUDA::cudart)
   endif()
-
+elseif(ENABLE_OMPGPU)
+  if(OMP_OFFLOAD_TARGET_NVIDIA GREATER 0)
+    target_compile_definitions(${FastCaloGpu_LIB} PUBLIC -DOMP_OFFLOAD_TARGET_NVIDIA)
+    target_link_libraries(${FastCaloGpu_LIB} PUBLIC ${CUDA_curand_LIBRARY} ${CUDA_nvToolsExt_LIBRARY})
+  endif()
+  if(OMP_OFFLOAD_TARGET_AMD GREATER 0)
+    target_compile_definitions(${FastCaloGpu_LIB} PUBLIC -D__HIP_PLATFORM_AMD__)
+    target_compile_definitions(${FastCaloGpu_LIB} PUBLIC -DOMP_OFFLOAD_TARGET_AMD)
+    find_package(HIP REQUIRED) 
+    target_include_directories(${FastCaloGpu_LIB} PRIVATE ${ROCM_PATH}/rocrand/include )
+    target_include_directories(${FastCaloGpu_LIB} PRIVATE ${ROCM_PATH}/include )
+    target_link_libraries(${FastCaloGpu_LIB}  PUBLIC ${ROCM_PATH}/lib/librocrand.so)
+    target_link_libraries(${FastCaloGpu_LIB}  PUBLIC ${ROCM_PATH}/lib/libamdhip64.so)
+  endif()
 else()
   target_link_libraries(${FastCaloGpu_LIB} PUBLIC ${CUDA_curand_LIBRARY} ${CUDA_nvToolsExt_LIBRARY})
 endif()
@@ -151,14 +179,58 @@ endif()
 if(RNDGEN_CPU)
   message(STATUS "Will generate random numbers on CPU")
   target_compile_definitions(${FastCaloGpu_LIB} PRIVATE -DRNDGEN_CPU )
-  # TODO Link a portable RNG library
-else()
-  if ( ${HIP_TARGET} STREQUAL "AMD" )
-    target_include_directories(${FastCaloGpu_LIB} PRIVATE ${ROCM_PATH}/hiprand/include )
+elseif(RNDGEN_OMP)
+  message(STATUS "Will generate random numbers using Portable OpenMP RNG Library")
+  target_compile_definitions(${FastCaloGpu_LIB} PRIVATE -DRNDGEN_OMP )
+  if(NOT DEFINED OMPRNG_HOME)
+    include(FetchContent)
+    set(FETCHCONTENT_QUIET OFF)
+    FetchContent_Declare(
+    Portable-OpenMP-RNG
+    GIT_REPOSITORY https://github.com/GKNB/test-benchmark-OpenMP-RNG
+    GIT_TAG        origin/main
+    )
+    FetchContent_Populate(Portable-OpenMP-RNG)
+    include_directories(${CMAKE_BINARY_DIR}/_deps/portable-openmp-rng-src)
+  else()
+    include_directories(${OMPRNG_HOME})
+    include_directories(${OMPRNG_HOME}/implementation)
+  endif()
+  if(ARCH_CUDA)
+    target_compile_definitions(${FastCaloGpu_LIB} PRIVATE -DARCH_CUDA )
+    target_link_libraries(${FastCaloGpu_LIB} PRIVATE ${CUDA_curand_LIBRARY} ${CUDA_nvToolsExt_LIBRARY})
+  elseif(ARCH_HIP)
+    target_compile_definitions(${FastCaloGpu_LIB} PRIVATE -DARCH_HIP )
     target_include_directories(${FastCaloGpu_LIB} PRIVATE ${ROCM_PATH}/rocrand/include )
-    target_link_libraries(${FastCaloGpu_LIB}  PUBLIC ${ROCM_PATH}/lib/libhiprand.so)
-  elseif( ${HIP_TARGET} STREQUAL "NVIDIA" )
-    target_link_libraries(${FastCaloGpu_LIB} PUBLIC ${CUDA_curand_LIBRARY} ${CUDA_nvToolsExt_LIBRARY})
+  elseif(USE_RANDOM123)
+    if(NOT DEFINED RANDOM123_HOME)
+      include(FetchContent)
+      set(FETCHCONTENT_QUIET OFF)
+      FetchContent_Declare(
+      random123
+      GIT_REPOSITORY https://github.com/DEShawResearch/random123
+      GIT_TAG        origin/main
+      )
+      FetchContent_Populate(random123)
+      include_directories(${CMAKE_BINARY_DIR}/_deps/random123-src/include )
+    else()
+      include_directories(${RANDOM123_HOME}/include/)
+    endif()
+    target_compile_definitions(${FastCaloGpu_LIB} PRIVATE -DUSE_RANDOM123 )
+  endif()
+else()
+  if(USE_HIP)
+    if ( ${HIP_TARGET} STREQUAL "AMD" )
+      target_include_directories(${FastCaloGpu_LIB} PRIVATE ${ROCM_PATH}/hiprand/include )
+      target_include_directories(${FastCaloGpu_LIB} PRIVATE ${ROCM_PATH}/rocrand/include )
+      target_link_libraries(${FastCaloGpu_LIB}  PUBLIC ${ROCM_PATH}/lib/libhiprand.so)
+    elseif( ${HIP_TARGET} STREQUAL "NVIDIA" )
+      target_link_libraries(${FastCaloGpu_LIB} PUBLIC ${CUDA_curand_LIBRARY} ${CUDA_nvToolsExt_LIBRARY})
+    endif()
+  elseif(ENABLE_OMPGPU)
+    if(ENV{OMP_TARGET_OFFLOAD} MATCHES "disabled")
+      message(FATAL_ERROR "when OMP_TARGET_OFFLOAD disabled, RNDGEN_CPU must be ON")
+    endif()
   endif()
 endif()
 

FastCaloSimAnalyzer/FastCaloGpu/src/CaloGpuGeneral_omp.cxx

@@ -8,16 +8,12 @@
 #include "Hit.h"
 #include "Rand4Hits.h"
 
-#include "gpuQ.h"
 #include "Args.h"
 #include "DEV_BigMem.h"
-// #include "OMP_BigMem.h"
 #include <chrono>
 #include <mutex>
 #include <climits>
 
-#include <cuda_runtime_api.h>
-#include <curand.h>
 #include <iostream>
 #include <omp.h>
 

FastCaloSimAnalyzer/FastCaloGpu/src/GeoRegion.cxx

@@ -6,13 +6,13 @@
 #include <iostream>
 #include <algorithm>
 
-#define PI 3.14159265358979323846
+#define PI_FCS 3.14159265358979323846
 #define TWOPI 2 * 3.14159265358979323846
 
 __HOSTDEV__ double Phi_mpi_pi(double x) {
-  while (x >= PI)
+  while (x >= PI_FCS)
     x -= TWOPI;
-  while (x < -PI)
+  while (x < -PI_FCS)
     x += TWOPI;
   return x;
 }

FastCaloSimAnalyzer/FastCaloGpu/src/Rand4Hits_omp.cxx

@@ -1,22 +1,35 @@
 /*
   Copyright (C) 2002-2021 CERN for the benefit of the ATLAS collaboration
 */
-#include "gpuQ.h"
 #include "Rand4Hits.h"
 #include "DEV_BigMem.h"
 
-#include <omp.h>
-#include <cuda_runtime_api.h>
-#include <curand.h>
-
 #include "GpuParams.h"
 #include "Rand4Hits_cpu.cxx"
 
-#define CURAND_CALL( x )                                                                                               \
-  if ( ( x ) != CURAND_STATUS_SUCCESS ) {                                                                              \
-    printf( "Error at %s:%d\n", __FILE__, __LINE__ );                                                                  \
-    exit( EXIT_FAILURE );                                                                                              \
-  }
+#include <omp.h>
+#ifdef RNDGEN_OMP
+#  include "openmp_rng.h"
+#endif
+
+#ifdef OMP_OFFLOAD_TARGET_NVIDIA
+#  include "gpuQ.h"
+#  include <cuda_runtime_api.h>
+#  include <curand.h>
+#  define CURAND_CALL( x )                                                                                               \
+    if ( ( x ) != CURAND_STATUS_SUCCESS ) {                                                                              \
+      printf( "Error at %s:%d\n", __FILE__, __LINE__ );                                                                  \
+      exit( EXIT_FAILURE );                                                                                              \
+    }
+#elif defined OMP_OFFLOAD_TARGET_AMD
+#  include "hip/hip_runtime.h"
+#  include <rocrand.h>
+#  define ROCRAND_CALL( x )                                                         \
+    if ((x) != ROCRAND_STATUS_SUCCESS) {                                          \
+      printf("Error at %s:%d\n", __FILE__, __LINE__);                            \
+      exit(EXIT_FAILURE);                                                        \
+    }
+#endif
 
 void Rand4Hits::allocate_simulation( int maxbins, int maxhitct, unsigned long n_cells ) {
 
@@ -67,12 +80,23 @@ Rand4Hits::~Rand4Hits() {
               << "  lost: " << DEV_BigMem::bm_ptr->lost() << std::endl;
     delete DEV_BigMem::bm_ptr;
   }
+#ifdef RNDGEN_OMP 	
   omp_target_free( m_rand_ptr, m_select_device );
+#endif
   if ( m_useCPU ) {
     destroyCPUGen();
   } else {
+#ifndef RNDGEN_OMP 	
+#ifndef USE_RANDOM123
+#ifdef OMP_OFFLOAD_TARGET_NVIDIA
     CURAND_CALL( curandDestroyGenerator( *( (curandGenerator_t*)m_gen ) ) );
     delete (curandGenerator_t*)m_gen;
+#elif defined OMP_OFFLOAD_TARGET_AMD
+    ROCRAND_CALL(rocrand_destroy_generator( *( (rocrand_generator*)m_gen)));
+    delete (rocrand_generator *)m_gen;
+#endif
+#endif
+#endif
   }
 };
 
@@ -84,7 +108,28 @@ void Rand4Hits::rd_regen() {
       std::cout << "ERROR: copy random numbers from cpu to gpu " << std::endl;
     }
   } else {
+#ifdef RNDGEN_OMP 	
+    auto gen = generator_enum::xorwow;
+#  ifdef USE_RANDOM123
+    float* f_r123 = (float*) malloc ( 3 * m_total_a_hits * sizeof( float ) );
+    omp_get_rng_uniform_float(f_r123, 3 * m_total_a_hits, m_seed, gen);
+    if ( omp_target_memcpy( m_rand_ptr, f_r123, 3 * m_total_a_hits * sizeof( float ), m_offset, m_offset, m_select_device,
+                            m_initial_device ) ) {
+      std::cout << "ERROR: copy random numbers from cpu to gpu " << std::endl;
+    }
+    free(f_r123);
+#  else
+    omp_get_rng_uniform_float(m_rand_ptr, 3 * m_total_a_hits, m_seed, gen);
+#  endif
+#else
+#  ifndef RNDGEN_CPU
+#  ifdef OMP_OFFLOAD_TARGET_NVIDIA
     CURAND_CALL( curandGenerateUniform( *( (curandGenerator_t*)m_gen ), m_rand_ptr, 3 * m_total_a_hits ) );
+#  elif defined OMP_OFFLOAD_TARGET_AMD
+    ROCRAND_CALL(rocrand_generate_uniform( *( (rocrand_generator*)m_gen), m_rand_ptr, 3 * m_total_a_hits));
+#  endif
+#  endif
+#endif 
   }
 };
 
@@ -105,12 +150,42 @@ void Rand4Hits::create_gen( unsigned long long seed, size_t num, bool useCPU ) {
       std::cout << "ERROR: copy random numbers from cpu to gpu " << std::endl;
     }
   } else {
+#ifdef RNDGEN_OMP 	
+    f = (float*)omp_target_alloc( num * sizeof( float ), m_select_device );
+    auto gen = generator_enum::xorwow; 
+ #ifdef USE_RANDOM123
+    float* f_r123 = (float*) malloc ( num * sizeof( float ) );
+    omp_get_rng_uniform_float(f_r123, num, seed, gen);
+    if ( omp_target_memcpy( f, f_r123, num * sizeof( float ), m_offset, m_offset, m_select_device,
+                            m_initial_device ) ) {
+      std::cout << "ERROR: copy random numbers from cpu to gpu " << std::endl;
+    }
+    free(f_r123);
+ #else
+    omp_get_rng_uniform_float(f, num, seed, gen);
+ #endif 
+    m_gen = (void*)gen;
+    // We need to save the seed for rd_regen
+    m_seed = seed;
+#else
+ #ifndef RNDGEN_CPU
+  #ifdef OMP_OFFLOAD_TARGET_NVIDIA
     gpuQ( cudaMalloc( &f, num * sizeof( float ) ) );
     curandGenerator_t* gen = new curandGenerator_t;
     CURAND_CALL( curandCreateGenerator( gen, CURAND_RNG_PSEUDO_DEFAULT ) );
     CURAND_CALL( curandSetPseudoRandomGeneratorSeed( *gen, seed ) );
     CURAND_CALL( curandGenerateUniform( *gen, f, num ) );
     m_gen = (void*)gen;
+  #elif defined OMP_OFFLOAD_TARGET_AMD
+    hipMalloc(&f, num * sizeof(float));
+    rocrand_generator* gen = new rocrand_generator;
+    ROCRAND_CALL(rocrand_create_generator(gen, ROCRAND_RNG_PSEUDO_DEFAULT));
+    ROCRAND_CALL(rocrand_set_seed(*gen, seed));
+    ROCRAND_CALL(rocrand_generate_uniform(*gen, f, num));
+    m_gen = (void*)gen;
+  #endif
+ #endif   
+#endif 
   }
 
   m_rand_ptr = f;

FastCaloSimAnalyzer/FastCaloGpu/src/gpuQ.cxx

@@ -3,6 +3,7 @@
 */
 
 #ifdef USE_OMPGPU
+#ifdef OMP_OFFLOAD_TARGET_NVIDIA
 #include "gpuQ.h"
 #include <iostream>
 
@@ -13,6 +14,7 @@ void gpu_assert(cudaError_t code, const char *file, const int line) {
     exit(code);
   }
 }
+#endif
 #else
 #include "gpuQ.cu"
 #endif