From effed42ef1d91ff156fa227f972aa7459eff14ba Mon Sep 17 00:00:00 2001 From: Willem Deconinck Date: Thu, 19 Sep 2024 10:59:08 +0300 Subject: [PATCH 1/3] Fix atlas_test_array when hic devices are present but no acc::devices --- src/tests/array/test_array.cc | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/src/tests/array/test_array.cc b/src/tests/array/test_array.cc index 290e47e0e..9f651fb13 100644 --- a/src/tests/array/test_array.cc +++ b/src/tests/array/test_array.cc @@ -20,6 +20,7 @@ #endif #include "hic/hic.h" +#include "atlas/parallel/acc/acc.h" using namespace atlas::array; @@ -555,7 +556,7 @@ CASE("test_wrap") { EXPECT(view(2) == 19); } -static int devices() { +static int hic_devices() { static int devices_ = [](){ int n = 0; auto err = hicGetDeviceCount(&n); @@ -573,7 +574,7 @@ CASE("test_acc_map") { EXPECT_NO_THROW(ds->allocateDevice()); if( ds->deviceAllocated() ) { EXPECT_NO_THROW(ds->accMap()); - EXPECT_EQ(ds->accMapped(), std::min(devices(),1)); + EXPECT_EQ(ds->accMapped(), std::min(std::min(acc::devices(),hic_devices()),1)); } else { Log::warning() << "WARNING: Array could not be allocated on device, so acc_map could not be tested" << std::endl; From 9af40a3312a11c9883fe167167fa2401a0149b28 Mon Sep 17 00:00:00 2001 From: Willem Deconinck Date: Thu, 19 Sep 2024 19:11:19 +0300 Subject: [PATCH 2/3] Update ACC support to use Fortran runtime functions, supporting Cray compilers --- cmake/features/ACC.cmake | 41 ++++------- cmake/project_summary.cmake | 4 +- src/atlas/array/gridtools/GridToolsArray.cc | 2 +- .../array/gridtools/GridToolsDataStore.h | 2 +- src/atlas/parallel/acc/acc.cc | 18 ++++- src/atlas_acc_support/CMakeLists.txt | 31 +++----- src/atlas_acc_support/atlas_acc.F90 | 70 +++++++++++++++++++ .../{atlas_acc_map_data.c => atlas_acc.cc} | 5 +- .../{atlas_acc_map_data.h => atlas_acc.h} | 12 +--- src/atlas_f/CMakeLists.txt | 12 ++-- src/sandbox/fortran_acc_fields/CMakeLists.txt | 4 +- src/tests/acc/CMakeLists.txt | 14 +--- .../fctest_unified_memory_with_openacc_cxx.cc | 2 +- src/tests/array/test_array.cc | 16 +---- src/tests/field/CMakeLists.txt | 16 ++--- src/tests/field/test_field_acc.cc | 6 +- 16 files changed, 137 insertions(+), 118 deletions(-) create mode 100644 src/atlas_acc_support/atlas_acc.F90 rename src/atlas_acc_support/{atlas_acc_map_data.c => atlas_acc.cc} (99%) rename src/atlas_acc_support/{atlas_acc_map_data.h => atlas_acc.h} (75%) diff --git a/cmake/features/ACC.cmake b/cmake/features/ACC.cmake index 056b2c812..273512006 100644 --- a/cmake/features/ACC.cmake +++ b/cmake/features/ACC.cmake @@ -1,40 +1,23 @@ ### OpenACC -if( atlas_HAVE_ATLAS_FIELD ) +if( atlas_HAVE_ATLAS_FIELD AND HAVE_GPU ) -set( ATLAS_ACC_CAPABLE FALSE ) -if( HAVE_GPU ) - if( CMAKE_Fortran_COMPILER_ID MATCHES "PGI|NVHPC" ) - set( ATLAS_ACC_CAPABLE TRUE ) - else() - find_package(OpenACC COMPONENTS C Fortran) - if(OpenACC_Fortran_FOUND AND OpenACC_C_FOUND) - set( ATLAS_ACC_CAPABLE TRUE ) - endif() + if( DEFINED ATLAS_ENABLE_ACC ) + set( ENABLE_ACC ${ATLAS_ENABLE_ACC} ) endif() -endif() - -ecbuild_add_option( FEATURE ACC - DESCRIPTION "OpenACC capable data structures" - CONDITION ATLAS_ACC_CAPABLE ) - -if( atlas_HAVE_ACC ) - if( CMAKE_Fortran_COMPILER_ID MATCHES "PGI|NVHPC" ) - #set( ACC_Fortran_FLAGS -acc -ta=tesla,nordc ) - set( ACC_Fortran_FLAGS "-acc=gpu;-gpu=gvmode,lineinfo,fastmath,rdc" ) - set( ACC_C_FLAGS ${ACC_Fortran_FLAGS} ) - find_program( ACC_C_COMPILER NAMES pgcc HINTS ${PGI_DIR} ${NVPHC_DIR} ENV PGI_DIR NVHPC_DIR PATH_SUFFIXES bin ) - if( NOT ACC_C_COMPILER ) - ecbuild_error( "Could not find OpenACC capable C compiler" ) + if( ENABLE_ACC ) + if( NOT HAVE_FORTRAN ) + enable_language(Fortran) endif() - else() - set( ACC_Fortran_FLAGS ${OpenACC_Fortran_FLAGS} ) - set( ACC_C_FLAGS ${OpenACC_C_FLAGS} ) + find_package( OpenACC COMPONENTS Fortran CXX ) endif() -endif() + ecbuild_add_option( FEATURE ACC + DESCRIPTION "OpenACC capable data structures" + CONDITION OpenACC_Fortran_FOUND ) else() + set( HAVE_ACC 0 ) set( atlas_HAVE_ACC 0 ) -endif() +endif() diff --git a/cmake/project_summary.cmake b/cmake/project_summary.cmake index 62ead6584..b5bfa6e3c 100644 --- a/cmake/project_summary.cmake +++ b/cmake/project_summary.cmake @@ -42,9 +42,7 @@ endif() if( atlas_HAVE_ACC ) ecbuild_info( "ACC" ) - ecbuild_info( " ACC_C_COMPILER : [${ACC_C_COMPILER}]" ) - ecbuild_info( " ACC_C_FLAGS : [${ACC_C_FLAGS}]" ) - ecbuild_info( " ACC_Fortran_FLAGS : [${ACC_Fortran_FLAGS}]" ) + ecbuild_info( " OpenACC_Fortran_FLAGS : [${OpenACC_Fortran_FLAGS}]" ) endif() diff --git a/src/atlas/array/gridtools/GridToolsArray.cc b/src/atlas/array/gridtools/GridToolsArray.cc index ac660742a..572c62385 100644 --- a/src/atlas/array/gridtools/GridToolsArray.cc +++ b/src/atlas/array/gridtools/GridToolsArray.cc @@ -25,7 +25,7 @@ #include "atlas/runtime/Log.h" #if ATLAS_HAVE_ACC -#include "atlas_acc_support/atlas_acc_map_data.h" +#include "atlas_acc_support/atlas_acc.h" #endif //------------------------------------------------------------------------------ diff --git a/src/atlas/array/gridtools/GridToolsDataStore.h b/src/atlas/array/gridtools/GridToolsDataStore.h index 8fb604640..058f726e5 100644 --- a/src/atlas/array/gridtools/GridToolsDataStore.h +++ b/src/atlas/array/gridtools/GridToolsDataStore.h @@ -15,7 +15,7 @@ #include "atlas/array/gridtools/GridToolsTraits.h" #if ATLAS_HAVE_ACC -#include "atlas_acc_support/atlas_acc_map_data.h" +#include "atlas_acc_support/atlas_acc.h" #endif //------------------------------------------------------------------------------ diff --git a/src/atlas/parallel/acc/acc.cc b/src/atlas/parallel/acc/acc.cc index 8a600aea5..c73bdea2c 100644 --- a/src/atlas/parallel/acc/acc.cc +++ b/src/atlas/parallel/acc/acc.cc @@ -13,7 +13,20 @@ #include "atlas/library/defines.h" #if ATLAS_HAVE_ACC -#include "atlas_acc_support/atlas_acc_map_data.h" +#include "hic/hic.h" +#include "atlas_acc_support/atlas_acc.h" +static int hic_devices() { + static int devices_ = [](){ + int n = 0; + auto err = hicGetDeviceCount(&n); + if (err != hicSuccess) { + n = 0; + static_cast(hicGetLastError()); + } + return n; + }(); + return devices_; +} #endif namespace atlas::acc { @@ -21,6 +34,9 @@ namespace atlas::acc { int devices() { #if ATLAS_HAVE_ACC static int num_devices = [](){ + if (hic_devices() == 0) { + return 0; + } auto devicetype = atlas_acc_get_device_type(); int _num_devices = atlas_acc_get_num_devices(); if (_num_devices == 1 && devicetype == atlas_acc_device_host) { diff --git a/src/atlas_acc_support/CMakeLists.txt b/src/atlas_acc_support/CMakeLists.txt index 7bcf8ac06..7bf818b1d 100644 --- a/src/atlas_acc_support/CMakeLists.txt +++ b/src/atlas_acc_support/CMakeLists.txt @@ -8,30 +8,15 @@ if( atlas_HAVE_ACC ) - if( NOT (CMAKE_C_COMPILER_ID MATCHES ${CMAKE_Fortran_COMPILER_ID}) ) - add_custom_command( - OUTPUT ${CMAKE_BINARY_DIR}/lib/libatlas_acc_support.so ${CMAKE_CURRENT_BINARY_DIR}/atlas_acc_map_data.c.o - COMMAND ${ACC_C_COMPILER} ${ACC_C_FLAGS} ${ACC_C_INCLUDE} -fPIC -o ${CMAKE_CURRENT_BINARY_DIR}/atlas_acc_map_data.c.o - -c ${CMAKE_CURRENT_SOURCE_DIR}/atlas_acc_map_data.c - COMMAND mkdir -p ${CMAKE_BINARY_DIR}/lib - COMMAND ${ACC_C_COMPILER} ${ACC_C_FLAGS} -shared -o ${CMAKE_BINARY_DIR}/lib/libatlas_acc_support.so - ${CMAKE_CURRENT_BINARY_DIR}/atlas_acc_map_data.c.o - DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/atlas_acc_map_data.c - COMMENT "Building atlas_acc_support with ${ACC_C_COMPILER}" - ) - add_custom_target( build-atlas_acc_support ALL DEPENDS ${CMAKE_BINARY_DIR}/lib/libatlas_acc_support.so ) - add_library( atlas_acc_support SHARED IMPORTED GLOBAL ) - set_property( TARGET atlas_acc_support PROPERTY IMPORTED_LOCATION ${CMAKE_BINARY_DIR}/lib/libatlas_acc_support.so ) - set_property( TARGET atlas_acc_support PROPERTY IMPORTED_NO_SONAME TRUE ) - set_property( TARGET atlas_acc_support PROPERTY IMPORTED_INCLUDE_DIRECTORIES ${CMAKE_CURRENT_SOURCE_DIR} ) - add_dependencies( atlas_acc_support build-atlas_acc_support ) - install( FILES ${CMAKE_BINARY_DIR}/lib/libatlas_acc_support.so DESTINATION ${INSTALL_LIB_DIR}/ ) - + if( CMAKE_CXX_COMPILER_ID MATCHES NVHPC ) + if( NOT TARGET OpenACC::OpenACC_CXX ) + ecbuild_error("ERROR: OpenACC::OpenACC_CXX TARGET not found") + endif() + ecbuild_add_library( TARGET atlas_acc_support SOURCES atlas_acc.cc ) + target_link_libraries( atlas_acc_support PRIVATE OpenACC::OpenACC_CXX ) else() - - ecbuild_add_library( TARGET atlas_acc_support SOURCES atlas_acc_map_data.c ) - target_compile_options( atlas_acc_support PRIVATE ${ACC_C_FLAGS} ) - target_link_libraries( atlas_acc_support PRIVATE ${ACC_C_FLAGS} ) + ecbuild_add_library( TARGET atlas_acc_support SOURCES atlas_acc.F90 ) + target_link_libraries( atlas_acc_support PRIVATE OpenACC::OpenACC_Fortran ) endif() endif() diff --git a/src/atlas_acc_support/atlas_acc.F90 b/src/atlas_acc_support/atlas_acc.F90 new file mode 100644 index 000000000..bb73b5237 --- /dev/null +++ b/src/atlas_acc_support/atlas_acc.F90 @@ -0,0 +1,70 @@ +module atlas_acc +use openacc +implicit none +private + +public :: atlas_acc_get_num_devices +public :: atlas_acc_map_data +public :: atlas_acc_unmap_data +public :: atlas_acc_is_present +public :: atlas_acc_get_device_type +public :: atlas_acc_deviceptr + +contains + +function atlas_acc_get_num_devices() bind(C,name="atlas_acc_get_num_devices") result(num_devices) + use, intrinsic :: iso_c_binding, only : c_int + integer(c_int) :: num_devices + integer(acc_device_kind) :: devicetype + + devicetype = acc_get_device_type() + num_devices = acc_get_num_devices(devicetype) +end function + +subroutine atlas_acc_map_data(data_arg, data_dev, bytes) bind(C,name="atlas_acc_map_data") + use, intrinsic :: iso_c_binding, only : c_ptr, c_size_t + type(*), dimension(*) :: data_arg + type(c_ptr), value :: data_dev + integer(c_size_t), value :: bytes + call acc_map_data(data_arg, data_dev, bytes) +end subroutine + +subroutine atlas_acc_unmap_data(data_arg) bind(C,name="atlas_acc_unmap_data") + use, intrinsic :: iso_c_binding, only : c_ptr + type(*), dimension(*) :: data_arg + call acc_unmap_data(data_arg) +end subroutine + +function atlas_acc_is_present(data_arg, bytes) bind(C,name="atlas_acc_is_present") result(is_present) + use, intrinsic :: iso_c_binding, only : c_size_t, c_ptr, c_char, c_int + integer(c_int) :: is_present + logical :: lpresent + type(c_ptr), value :: data_arg + integer(c_size_t), value :: bytes + character(kind=c_char), pointer :: data_f(:) + call c_f_pointer(data_arg, data_f,[bytes]) + lpresent = acc_is_present(data_f) + is_present = 0 + if (lpresent) is_present = 1 +end function + +function atlas_acc_deviceptr(data_arg) bind(C,name="atlas_acc_deviceptr") result(deviceptr) + use, intrinsic :: iso_c_binding, only : c_ptr + type(*), dimension(*) :: data_arg + type(c_ptr):: deviceptr + deviceptr = acc_deviceptr(data_arg) +end function + +function atlas_acc_get_device_type() bind(C,name="atlas_acc_get_device_type") result(devicetype) + use, intrinsic :: iso_c_binding, only : c_int + integer(c_int) :: devicetype + integer(acc_device_kind) :: acc_devicetype + acc_devicetype = acc_get_device_type() + if (acc_devicetype == acc_device_host .or. acc_devicetype == acc_device_none) then + devicetype = 0 + else + devicetype = 1 + endif +end function + +end module diff --git a/src/atlas_acc_support/atlas_acc_map_data.c b/src/atlas_acc_support/atlas_acc.cc similarity index 99% rename from src/atlas_acc_support/atlas_acc_map_data.c rename to src/atlas_acc_support/atlas_acc.cc index 826579686..8cc58981a 100644 --- a/src/atlas_acc_support/atlas_acc_map_data.c +++ b/src/atlas_acc_support/atlas_acc.cc @@ -19,7 +19,9 @@ #include #include -#include "atlas_acc_map_data.h" +#include "atlas_acc.h" + +extern "C" { void atlas_acc_map_data(void* cpu_ptr, void* gpu_ptr, unsigned long bytes) { acc_map_data(cpu_ptr, gpu_ptr, bytes); @@ -125,3 +127,4 @@ const char* atlas_acc_info_str() { int atlas_acc_get_num_devices() { return acc_get_num_devices(acc_get_device_type()); } +} diff --git a/src/atlas_acc_support/atlas_acc_map_data.h b/src/atlas_acc_support/atlas_acc.h similarity index 75% rename from src/atlas_acc_support/atlas_acc_map_data.h rename to src/atlas_acc_support/atlas_acc.h index ba5990968..37c928b95 100644 --- a/src/atlas_acc_support/atlas_acc_map_data.h +++ b/src/atlas_acc_support/atlas_acc.h @@ -10,6 +10,7 @@ #pragma once +#include #ifdef __cplusplus extern "C" { #endif @@ -19,18 +20,11 @@ typedef enum { atlas_acc_device_not_host = 1 } atlas_acc_device_t; -void atlas_acc_map_data(void* cpu_ptr, void* gpu_ptr, unsigned long bytes); +void atlas_acc_map_data(void* cpu_ptr, void* gpu_ptr, size_t bytes); void atlas_acc_unmap_data(void* cpu_ptr); -int atlas_acc_is_present(void* cpu_ptr, unsigned long bytes); +int atlas_acc_is_present(void* cpu_ptr, size_t bytes); void* atlas_acc_deviceptr(void* cpu_ptr); atlas_acc_device_t atlas_acc_get_device_type(); - -int atlas_acc_devices(); - -const char* atlas_acc_version_str(); - -const char* atlas_acc_info_str(); - int atlas_acc_get_num_devices(); #ifdef __cplusplus diff --git a/src/atlas_f/CMakeLists.txt b/src/atlas_f/CMakeLists.txt index 2e9a97e31..53067f222 100644 --- a/src/atlas_f/CMakeLists.txt +++ b/src/atlas_f/CMakeLists.txt @@ -278,11 +278,13 @@ ecbuild_add_library( TARGET atlas_f ) if( HAVE_ACC ) - target_link_options( atlas_f INTERFACE - $<$:SHELL:-acc=gpu> - $<$:SHELL:-acc=gpu> - $<$:SHELL:-acc=gpu> - $<$:SHELL:-acc=gpu> ) + if( CMAKE_Fortran_COMPILER_ID MATCHES NVHPC ) + target_link_options( atlas_f INTERFACE + $<$:SHELL:-acc=gpu> + $<$:SHELL:-acc=gpu> + $<$:SHELL:-acc=gpu> + $<$:SHELL:-acc=gpu> ) + endif() endif() fckit_target_preprocess_fypp( atlas_f diff --git a/src/sandbox/fortran_acc_fields/CMakeLists.txt b/src/sandbox/fortran_acc_fields/CMakeLists.txt index a3b49ad08..733960140 100644 --- a/src/sandbox/fortran_acc_fields/CMakeLists.txt +++ b/src/sandbox/fortran_acc_fields/CMakeLists.txt @@ -10,10 +10,8 @@ if( atlas_HAVE_ACC ) ecbuild_add_executable( TARGET atlas-acc-fields SOURCES atlas-acc-fields.F90 - LIBS atlas_f + LIBS atlas_f OpenACC::OpenACC_Fortran LINKER_LANGUAGE Fortran NOINSTALL ) - target_compile_options( atlas-acc-fields PUBLIC ${ACC_Fortran_FLAGS} ) - target_link_libraries( atlas-acc-fields ${ACC_Fortran_FLAGS} ) endif() diff --git a/src/tests/acc/CMakeLists.txt b/src/tests/acc/CMakeLists.txt index 0e323611f..c827dc65a 100644 --- a/src/tests/acc/CMakeLists.txt +++ b/src/tests/acc/CMakeLists.txt @@ -6,33 +6,25 @@ # granted to it by virtue of its status as an intergovernmental organisation nor # does it submit to any jurisdiction. -if( HAVE_CUDA AND HAVE_TESTS AND HAVE_FCTEST AND HAVE_ACC ) +if( HAVE_GPU AND HAVE_TESTS AND HAVE_FCTEST AND HAVE_ACC ) - string (REPLACE ";" " " ACC_Fortran_FLAGS_STR "${ACC_Fortran_FLAGS}") - - - set_source_files_properties( fctest_unified_memory_with_openacc.F90 PROPERTIES COMPILE_FLAGS ${ACC_Fortran_FLAGS_STR} ) add_fctest( TARGET atlas_test_unified_memory_with_openacc SOURCES fctest_unified_memory_with_openacc.F90 fctest_unified_memory_with_openacc_cxx.cc - LIBS atlas_f + LIBS atlas_f OpenACC::OpenACC_Fortran hic LINKER_LANGUAGE Fortran ENVIRONMENT ${ATLAS_TEST_ENVIRONMENT} ATLAS_RUN_NGPUS=1 ) - target_link_libraries( atlas_test_unified_memory_with_openacc ${ACC_Fortran_FLAGS} hic ) - set_tests_properties( atlas_test_unified_memory_with_openacc PROPERTIES LABELS "gpu;acc") add_fctest( TARGET atlas_test_connectivity_openacc SOURCES fctest_connectivity_openacc.F90 - LIBS atlas_f + LIBS atlas_f OpenACC::OpenACC_Fortran LINKER_LANGUAGE Fortran ENVIRONMENT ${ATLAS_TEST_ENVIRONMENT} ATLAS_RUN_NGPUS=1 ) - target_link_libraries( atlas_test_connectivity_openacc ${ACC_Fortran_FLAGS} ) - set_target_properties( atlas_test_connectivity_openacc PROPERTIES COMPILE_FLAGS "${ACC_Fortran_FLAGS_STR}" ) set_tests_properties ( atlas_test_connectivity_openacc PROPERTIES LABELS "gpu;acc") endif() diff --git a/src/tests/acc/fctest_unified_memory_with_openacc_cxx.cc b/src/tests/acc/fctest_unified_memory_with_openacc_cxx.cc index eb251572b..4ded2d78e 100644 --- a/src/tests/acc/fctest_unified_memory_with_openacc_cxx.cc +++ b/src/tests/acc/fctest_unified_memory_with_openacc_cxx.cc @@ -12,6 +12,6 @@ extern "C" { void allocate_unified_impl(double** a, int N) { - hicMallocManaged(a, N * sizeof(double)); + HIC_CALL(hicMallocManaged(a, N * sizeof(double))); } } diff --git a/src/tests/array/test_array.cc b/src/tests/array/test_array.cc index 9f651fb13..63f630c98 100644 --- a/src/tests/array/test_array.cc +++ b/src/tests/array/test_array.cc @@ -19,7 +19,6 @@ #include "atlas/array/gridtools/GridToolsMakeView.h" #endif -#include "hic/hic.h" #include "atlas/parallel/acc/acc.h" using namespace atlas::array; @@ -556,25 +555,12 @@ CASE("test_wrap") { EXPECT(view(2) == 19); } -static int hic_devices() { - static int devices_ = [](){ - int n = 0; - auto err = hicGetDeviceCount(&n); - if (err != hicSuccess) { - n = 0; - static_cast(hicGetLastError()); - } - return n; - }(); - return devices_; -} - CASE("test_acc_map") { Array* ds = Array::create(2, 3, 4); EXPECT_NO_THROW(ds->allocateDevice()); if( ds->deviceAllocated() ) { EXPECT_NO_THROW(ds->accMap()); - EXPECT_EQ(ds->accMapped(), std::min(std::min(acc::devices(),hic_devices()),1)); + EXPECT_EQ(ds->accMapped(), std::min(acc::devices(),1)); } else { Log::warning() << "WARNING: Array could not be allocated on device, so acc_map could not be tested" << std::endl; diff --git a/src/tests/field/CMakeLists.txt b/src/tests/field/CMakeLists.txt index 3382e576b..dc91d6f65 100644 --- a/src/tests/field/CMakeLists.txt +++ b/src/tests/field/CMakeLists.txt @@ -27,13 +27,11 @@ ecbuild_add_test( TARGET atlas_test_field_foreach ecbuild_add_test( TARGET atlas_test_field_acc SOURCES test_field_acc.cc - LIBS atlas + LIBS atlas OpenACC::OpenACC_CXX ENVIRONMENT ${ATLAS_TEST_ENVIRONMENT} - CONDITION atlas_HAVE_ACC + CONDITION atlas_HAVE_ACC AND OpenACC_CXX_Found ) if( TEST atlas_test_field_acc ) - target_compile_options( atlas_test_field_acc PRIVATE "${ACC_C_FLAGS}") - target_link_options( atlas_test_field_acc PRIVATE "${ACC_C_FLAGS}") set_tests_properties ( atlas_test_field_acc PROPERTIES LABELS "gpu;acc") endif() @@ -73,14 +71,11 @@ if( HAVE_FCTEST ) CONDITION atlas_HAVE_ACC LINKER_LANGUAGE Fortran SOURCES fctest_field_gpu.F90 external_acc_routine.F90 - LIBS atlas_f + LIBS atlas_f OpenACC::OpenACC_Fortran ENVIRONMENT ${ATLAS_TEST_ENVIRONMENT} ATLAS_RUN_NGPUS=1 ) if( TARGET atlas_fctest_field_device ) - target_compile_options( atlas_fctest_field_device PUBLIC ${ACC_Fortran_FLAGS} ) - target_link_libraries( atlas_fctest_field_device ${ACC_Fortran_FLAGS} ) - target_link_options( atlas_fctest_field_device PRIVATE "${ACC_Fortran_FLAGS}") set_tests_properties ( atlas_fctest_field_device PROPERTIES LABELS "gpu;acc") endif() @@ -88,14 +83,11 @@ if( HAVE_FCTEST ) CONDITION atlas_HAVE_ACC LINKER_LANGUAGE Fortran SOURCES fctest_field_wrap_gpu.F90 external_acc_routine.F90 - LIBS atlas_f + LIBS atlas_f OpenACC::OpenACC_Fortran ENVIRONMENT ${ATLAS_TEST_ENVIRONMENT} ATLAS_RUN_NGPUS=1 ) if( TARGET atlas_fctest_field_wrap_device ) - target_compile_options( atlas_fctest_field_wrap_device PUBLIC ${ACC_Fortran_FLAGS} ) - target_link_libraries( atlas_fctest_field_wrap_device ${ACC_Fortran_FLAGS} ) - target_link_options( atlas_fctest_field_wrap_device PRIVATE "${ACC_Fortran_FLAGS}") set_tests_properties ( atlas_fctest_field_wrap_device PROPERTIES LABELS "gpu;acc") endif() diff --git a/src/tests/field/test_field_acc.cc b/src/tests/field/test_field_acc.cc index e998f5dda..8dad81fbd 100644 --- a/src/tests/field/test_field_acc.cc +++ b/src/tests/field/test_field_acc.cc @@ -35,10 +35,10 @@ CASE("test_acc") { *c_ptr = 5; int* d_ptr; - hicMalloc(&d_ptr, sizeof(int)); + HIC_CALL(hicMalloc(&d_ptr, sizeof(int))); acc_map_data(c_ptr, d_ptr, sizeof(int)); - hicMemcpy(d_ptr, c_ptr, sizeof(int), hicMemcpyHostToDevice); + HIC_CALL(hicMemcpy(d_ptr, c_ptr, sizeof(int), hicMemcpyHostToDevice)); #pragma acc kernels present(c_ptr) { @@ -47,7 +47,7 @@ CASE("test_acc") { EXPECT_EQ( *c_ptr, 5. ); - hicMemcpy(c_ptr, d_ptr, sizeof(int), hicMemcpyDeviceToHost); + HIC_CALL(hicMemcpy(c_ptr, d_ptr, sizeof(int), hicMemcpyDeviceToHost)); EXPECT_EQ( *c_ptr, 2. ); } From cb0270c9cb50cd43faa038d04e9ff25637edb79a Mon Sep 17 00:00:00 2001 From: Willem Deconinck Date: Tue, 24 Sep 2024 15:46:17 +0000 Subject: [PATCH 3/3] Make OpenACC link options cleaner --- cmake/features/ACC.cmake | 4 ++++ src/atlas/CMakeLists.txt | 10 +++++----- src/atlas_f/CMakeLists.txt | 14 ++++++-------- 3 files changed, 15 insertions(+), 13 deletions(-) diff --git a/cmake/features/ACC.cmake b/cmake/features/ACC.cmake index 273512006..a1537142e 100644 --- a/cmake/features/ACC.cmake +++ b/cmake/features/ACC.cmake @@ -15,6 +15,10 @@ if( atlas_HAVE_ATLAS_FIELD AND HAVE_GPU ) ecbuild_add_option( FEATURE ACC DESCRIPTION "OpenACC capable data structures" CONDITION OpenACC_Fortran_FOUND ) + if( HAVE_ACC ) + set( ACC_LINK_OPTIONS ${OpenACC_Fortran_FLAGS} ) + endif() + else() set( HAVE_ACC 0 ) diff --git a/src/atlas/CMakeLists.txt b/src/atlas/CMakeLists.txt index 47c82b031..50091b3e8 100644 --- a/src/atlas/CMakeLists.txt +++ b/src/atlas/CMakeLists.txt @@ -1001,12 +1001,12 @@ ecbuild_add_library( TARGET atlas ) -if( HAVE_ACC ) +if( HAVE_ACC AND CMAKE_Fortran_COMPILER_ID MATCHES NVHPC ) target_link_options( atlas INTERFACE - $<$:SHELL:-acc=gpu> - $<$:SHELL:-acc=gpu> - $<$:SHELL:-acc=gpu> - $<$:SHELL:-acc=gpu> ) + $<$:SHELL:${ACC_LINK_OPTIONS}> + $<$:SHELL:${ACC_LINK_OPTIONS}> + $<$:SHELL:${ACC_LINK_OPTIONS}> + $<$:SHELL:${ACC_LINK_OPTIONS}> ) endif() target_compile_features( atlas PUBLIC cxx_std_17 ) diff --git a/src/atlas_f/CMakeLists.txt b/src/atlas_f/CMakeLists.txt index 53067f222..4fc068f9c 100644 --- a/src/atlas_f/CMakeLists.txt +++ b/src/atlas_f/CMakeLists.txt @@ -277,14 +277,12 @@ ecbuild_add_library( TARGET atlas_f $ ) -if( HAVE_ACC ) - if( CMAKE_Fortran_COMPILER_ID MATCHES NVHPC ) - target_link_options( atlas_f INTERFACE - $<$:SHELL:-acc=gpu> - $<$:SHELL:-acc=gpu> - $<$:SHELL:-acc=gpu> - $<$:SHELL:-acc=gpu> ) - endif() +if( HAVE_ACC AND CMAKE_Fortran_COMPILER_ID MATCHES NVHPC ) + target_link_options( atlas_f INTERFACE + $<$:SHELL:${ACC_LINK_OPTIONS}> + $<$:SHELL:${ACC_LINK_OPTIONS}> + $<$:SHELL:${ACC_LINK_OPTIONS}> + $<$:SHELL:${ACC_LINK_OPTIONS}> ) endif() fckit_target_preprocess_fypp( atlas_f