Add cuda tests (#18)

lukasm91 · havogt · commit a4281f956f7a · 2019-07-02T09:07:51.000+02:00
Adds a specific test for OpenACC support in array wrappers. Tests only compilation, but not execution because obviously, it's nearly impossible to have only one toolchain that compiles fortran and C++ successfully. Enabling the test for execution should be revisited as gcc gets better in OpenACC support.
diff --git a/cmake/cpp_bindgen.cmake.in b/cmake/cpp_bindgen.cmake.in
@@ -110,7 +110,6 @@ function(cpp_bindgen_add_library target_name)
 
     add_library(${target_name} ${ARG_SOURCES})
     target_link_libraries(${target_name} PRIVATE cpp_bindgen_generator)
-    # target_include_directories(${target_name} PRIVATE ${PROJECT_SOURCE_DIR}/include) #TODO probably wrong
 
     if(GT_ENABLE_BINDINGS_GENERATION)
         # generator
diff --git a/cmake/detect_features.cmake b/cmake/detect_features.cmake
@@ -1,11 +1,11 @@
 if(NOT DEFINED CPP_BINDGEN_ENABLE_COMPILER_DETECTION)
     if(${PROJECT_SOURCE_DIR} STREQUAL ${CMAKE_SOURCE_DIR})
+        set(CPP_BINDGEN_ENABLE_COMPILER_DETECTION ON)
+    else()
         # If turned OFF, it will still use the compilers,
         # - if they are available from a super-project, or
         # - if compilers are forced to on via other options.
         set(CPP_BINDGEN_ENABLE_COMPILER_DETECTION OFF)
-    else()
-        set(CPP_BINDGEN_ENABLE_COMPILER_DETECTION ON)
     endif()
 endif()
 
diff --git a/example/simple_fetch_content/CMakeLists.txt b/example/simple_fetch_content/CMakeLists.txt
@@ -1,21 +1,22 @@
 cmake_minimum_required(VERSION 3.12.4) #TODO replace with generated version
 project(simple_example_with_fetch_content LANGUAGES CXX Fortran)
 
-# 1) fetch c_bindings from the repository
+# 1) fetch cpp_bindgen from the repository
 # TODO: use FetchContent_MakeAvailable() once we require CMake 3.14+
 include(FetchContent)
 FetchContent_Declare(
-  c_bindings
+  cpp_bindgen
   # In your application remove SOURCE_DIR ...
   SOURCE_DIR       ${CMAKE_CURRENT_LIST_DIR}/../..
   # By the following 2 lines to fetch content directly from github
-  # GIT_REPOSITORY https://github.com/GridTools/c_bindings.git
+  # GIT_REPOSITORY https://github.com/GridTools/cpp_bindgen.git
   # GIT_TAG        master # consider replacing master by a tagged version
 )
-FetchContent_GetProperties(c_bindings)
+# FetchContent_MakeAvailable(cpp_bindgen) for CMake >= 3.14
+FetchContent_GetProperties(cpp_bindgen)
 if(NOT c_bindings_POPULATED)
-  FetchContent_Populate(c_bindings)
-  add_subdirectory(${c_bindings_SOURCE_DIR} ${c_bindings_BINARY_DIR})
+  FetchContent_Populate(cpp_bindgen)
+  add_subdirectory(${cpp_bindgen_SOURCE_DIR} ${cpp_bindgen_BINARY_DIR})
 endif()
 
 # 2) create a library with bindings. This will generate the files simple.h and simple.f90 which can be included within C and Fortran. In CMake you can use them by linking against `simple_c`or `simple_fortran`.
diff --git a/jenkins/build.sh b/jenkins/build.sh
@@ -3,6 +3,7 @@
 set -e
 
 HOST=$(hostname)
+test -n "${CSCS_ACCOUNT}" || CSCS_ACCOUNT=d75
 
 if [[ "$HOST" == kesch* || "$HOST" == escha* ]]; then
     module load cmake/3.13.4
@@ -15,11 +16,15 @@ if [[ "$HOST" == kesch* || "$HOST" == escha* ]]; then
     export CUDATOOLKIT_HOME=$CUDA_PATH
     export CUDA_ARCH=sm_37
     export FC=`which gfortran`
+    RUN_PREFIX="srun -p pp-short -c 12 --time=00:30:00"
+
 elif [[ "$HOST" == tave* ]]; then
     module switch PrgEnv-cray PrgEnv-gnu
     module rm CMake
     module load /users/jenkins/easybuild/tave/modules/all/CMake/3.12.4
     export BOOST_ROOT=/project/c14/install/kesch/boost/boost_1_67_0
+    RUN_PREFIX=""
+
 elif [[ "$HOST" == daint* ]]; then
     module load daint-gpu
     module load cudatoolkit/9.2.148_3.19-6.0.7.1_2.1__g3d9acc8
@@ -28,6 +33,8 @@ elif [[ "$HOST" == daint* ]]; then
     module load /users/jenkins/easybuild/daint/haswell/modules/all/CMake/3.12.4
     export BOOST_ROOT=/project/c14/install/daint/boost/boost_1_67_0
     export CUDA_ARCH=sm_60
+    RUN_PREFIX="srun -C gpu -p cscsci --account=$CSCS_ACCOUNT --time=00:30:00"
+
 else
     echo "Unknown host ${HOST}, using current environment."
 fi
@@ -38,7 +45,7 @@ cwd=$(pwd)
 mkdir -p build && cd build
 cmake .. -DCMAKE_INSTALL_PREFIX=${cwd}/install -DCPP_BINDGEN_REQUIRE_TEST_C=ON -DCPP_BINDGEN_REQUIRE_TEST_Fortran=ON
 nice make -j8 install
-ctest .
+$RUN_PREFIX ctest .
 
 # compile example using the config from build tree (test for export(package))
 cd ${cwd}/example/simple
@@ -68,5 +75,5 @@ if [[ "$HOST" != kesch* ]]; then
     mkdir -p build_legacy && cd build_legacy
     cmake .. -DCMAKE_INSTALL_PREFIX=${cwd}/install -DCPP_BINDGEN_GT_LEGACY=ON
     nice make -j8 install
-    ctest .
+    $RUN_PREFIX ctest .
 fi
diff --git a/tests/regression/array/CMakeLists.txt b/tests/regression/array/CMakeLists.txt
@@ -4,4 +4,17 @@ add_executable(gen_regression_array_driver_fortran driver.f90)
 target_link_libraries(gen_regression_array_driver_fortran gen_regression_array_fortran)
 add_test(NAME gen_regression_array_driver_fortran
     COMMAND $<TARGET_FILE:gen_regression_array_driver_fortran>)
-set_target_properties(gen_regression_array_driver_fortran PROPERTIES LINKER_LANGUAGE Fortran)
+
+if(CUDA_AVAILABLE AND "${CMAKE_CXX_COMPILER_VERSION}" VERSION_GREATER_EQUAL "6")
+    cpp_bindgen_add_library(gen_regression_array_cu SOURCES implementation.cu)
+    cpp_bindgen_enable_fortran_openacc_on_target(gen_regression_array_cu_fortran)
+
+    add_executable(gen_regression_array_driver_cu_fortran driver_cu.f90)
+    target_link_libraries(gen_regression_array_driver_cu_fortran gen_regression_array_cu_fortran)
+    cpp_bindgen_enable_fortran_openacc_on_target(gen_regression_array_driver_cu_fortran)
+
+    # TODO
+    # we don't add this test on purpose for now, because it is likely that this test does not work
+    # out of the box. The Fortran code would need to be compiled separately with a different
+    # compiler than the C++ code with a compiler that has proper OpenACC support.
+endif()
diff --git a/tests/regression/array/driver.f90 b/tests/regression/array/driver.f90
@@ -17,13 +17,15 @@ program main
 
     call fill_array(arr)
 
-    DO i=1, ie
-        DO j=1, je
-            DO k=1, ke
+    do i=1, ie
+        do j=1, je
+            do k=1, ke
                 expected(i,j,k) = (i-1)*10000 + (j-1)*100 + (k-1)
-            END DO
-        END DO
-    END DO
+            end do
+        end do
+    end do
 
     if (any(arr /= expected)) stop 1
+
+    print *, "Success!"
 end
diff --git a/tests/regression/array/driver_cu.f90 b/tests/regression/array/driver_cu.f90
@@ -0,0 +1,35 @@
+! GridTools
+!
+! Copyright (c) 2014-2019, ETH Zurich
+! All rights reserved.
+!
+! Please, refer to the LICENSE file in the root directory.
+! SPDX-License-Identifier: BSD-3-Clause
+
+program main
+    use iso_c_binding
+    use gen_handle
+    use gen_regression_array_cu
+    implicit none
+    integer, parameter :: ie = 9, je = 10, ke = 11
+    integer :: i, j, k
+    real(8), dimension(ie, je, ke) :: arr, expected
+
+    !$acc enter data create(arr)
+
+    call fill_gpu_array(arr)
+
+    !$acc exit data copyout(arr)
+
+    do i=1, ie
+        do j=1, je
+            do k=1, ke
+                expected(i,j,k) = (i-1)*10000 + (j-1)*100 + (k-1)
+            end do
+        end do
+    end do
+
+    if (any(arr /= expected)) stop 1
+
+    print *, "Success!"
+end
diff --git a/tests/regression/array/gen_regression_array_cu.f90 b/tests/regression/array/gen_regression_array_cu.f90
@@ -0,0 +1,51 @@
+! This file is generated!
+module gen_regression_array_cu
+implicit none
+  interface
+
+    subroutine fill_array_impl(arg0) bind(c, name="fill_array")
+      use iso_c_binding
+      use gen_array_descriptor
+      type(gen_fortran_array_descriptor) :: arg0
+    end subroutine
+    subroutine fill_gpu_array_impl(arg0) bind(c, name="fill_gpu_array")
+      use iso_c_binding
+      use gen_array_descriptor
+      type(gen_fortran_array_descriptor) :: arg0
+    end subroutine
+
+  end interface
+contains
+    subroutine fill_array(arg0)
+      use iso_c_binding
+      use gen_array_descriptor
+      real(c_double), dimension(:,:,:), target :: arg0
+      type(gen_fortran_array_descriptor) :: descriptor0
+
+      descriptor0%rank = 3
+      descriptor0%type = 6
+      descriptor0%dims = reshape(shape(arg0), &
+        shape(descriptor0%dims), (/0/))
+      descriptor0%data = c_loc(arg0(lbound(arg0, 1),lbound(arg0, 2),lbound(arg0, 3)))
+
+      call fill_array_impl(descriptor0)
+    end subroutine
+    subroutine fill_gpu_array(arg0)
+      use iso_c_binding
+      use gen_array_descriptor
+      real(c_double), dimension(:,:,:), target :: arg0
+      type(gen_fortran_array_descriptor) :: descriptor0
+
+      !$acc data present(arg0)
+      !$acc host_data use_device(arg0)
+      descriptor0%rank = 3
+      descriptor0%type = 6
+      descriptor0%dims = reshape(shape(arg0), &
+        shape(descriptor0%dims), (/0/))
+      descriptor0%data = c_loc(arg0(lbound(arg0, 1),lbound(arg0, 2),lbound(arg0, 3)))
+      !$acc end host_data
+      !$acc end data
+
+      call fill_gpu_array_impl(descriptor0)
+    end subroutine
+end
diff --git a/tests/regression/array/gen_regression_array_cu.h b/tests/regression/array/gen_regression_array_cu.h
@@ -0,0 +1,16 @@
+// This file is generated!
+#pragma once
+
+#include <cpp_bindgen/array_descriptor.h>
+#include <cpp_bindgen/handle.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+void fill_array(gen_fortran_array_descriptor*);
+void fill_gpu_array(gen_fortran_array_descriptor*);
+
+#ifdef __cplusplus
+}
+#endif
diff --git a/tests/regression/array/implementation.cpp b/tests/regression/array/implementation.cpp
@@ -11,7 +11,7 @@
 #include <array>
 
 #include <cpp_bindgen/export.hpp>
-#include <type_traits>
+#include <iostream>
 
 namespace custom_array {
     template <class T>
diff --git a/tests/regression/array/implementation.cu b/tests/regression/array/implementation.cu
@@ -0,0 +1,72 @@
+/*
+ * GridTools
+ *
+ * Copyright (c) 2014-2019, ETH Zurich
+ * All rights reserved.
+ *
+ * Please, refer to the LICENSE file in the root directory.
+ * SPDX-License-Identifier: BSD-3-Clause
+ */
+#include "implementation.cpp"
+
+#include <iostream>
+
+namespace gpu_array {
+    template <class T>
+    struct my_array {
+        using data_t = T;
+
+        T *data;
+        int sizes[3];
+        int strides[3];
+
+        __device__ const T &operator()(int i, int j, int k) const {
+            assert(i < sizes[0] && j < sizes[1] && k < sizes[2] && "out of bounds");
+            return data[i * strides[0] + j * strides[1] + k * strides[2]];
+        }
+
+        __device__ T &operator()(int i, int j, int k) {
+            assert(i < sizes[0] && j < sizes[1] && k < sizes[2] && "out of bounds");
+            return data[i * strides[0] + j * strides[1] + k * strides[2]];
+        }
+    };
+
+    template <typename T>
+    my_array<T> gen_make_fortran_array_view(gen_fortran_array_descriptor *descriptor, my_array<T> *) {
+        if (descriptor->rank != 3) {
+            throw std::runtime_error("only 3-dimensional arrays are supported");
+        }
+        cudaPointerAttributes attributes;
+        auto ret = cudaPointerGetAttributes(&attributes, descriptor->data);
+        if (ret != cudaSuccess || attributes.memoryType != cudaMemoryTypeDevice) {
+            throw std::runtime_error("no gpu pointer");
+        }
+        return my_array<T>{static_cast<T *>(descriptor->data),
+            {descriptor->dims[0], descriptor->dims[1], descriptor->dims[2]},
+            {1, descriptor->dims[0], descriptor->dims[0] * descriptor->dims[1]}};
+    }
+
+    template <typename T>
+    gen_fortran_array_descriptor get_fortran_view_meta(my_array<T> *) {
+        gen_fortran_array_descriptor descriptor;
+        descriptor.type = cpp_bindgen::fortran_array_element_kind<T>::value;
+        descriptor.rank = 3;
+        descriptor.is_acc_present = true;
+        return descriptor;
+    }
+
+    static_assert(cpp_bindgen::is_fortran_array_bindable<my_array<double>>::value, "");
+    static_assert(cpp_bindgen::is_fortran_array_wrappable<my_array<double>>::value, "");
+} // namespace gpu_array
+
+namespace {
+    __global__ void fill_array_kernel(gpu_array::my_array<double> a) {
+        for (size_t i = 0; i < a.sizes[2]; ++i) {
+            a(threadIdx.x, blockIdx.x, i) = threadIdx.x * 10000 + blockIdx.x * 100 + i;
+        }
+    }
+
+    void fill_gpu_array_impl(gpu_array::my_array<double> a) { fill_array_kernel<<<a.sizes[1], a.sizes[0]>>>(a); }
+
+    GEN_EXPORT_BINDING_WRAPPED_1(fill_gpu_array, fill_gpu_array_impl);
+} // namespace
diff --git a/tests/regression/array_gt_legacy/CMakeLists.txt b/tests/regression/array_gt_legacy/CMakeLists.txt
@@ -4,4 +4,3 @@ add_executable(gen_regression_array_gt_legacy_driver_fortran driver.f90)
 target_link_libraries(gen_regression_array_gt_legacy_driver_fortran gen_regression_array_gt_legacy_fortran)
 add_test(NAME gen_regression_array_gt_legacy_driver_fortran
     COMMAND $<TARGET_FILE:gen_regression_array_gt_legacy_driver_fortran>)
-set_target_properties(gen_regression_array_gt_legacy_driver_fortran PROPERTIES LINKER_LANGUAGE Fortran)
diff --git a/tests/regression/simple/CMakeLists.txt b/tests/regression/simple/CMakeLists.txt
@@ -5,7 +5,6 @@ if(CMAKE_Fortran_COMPILER_LOADED)
     target_link_libraries(gen_regression_simple_driver_fortran gen_regression_simple_fortran)
     add_test(NAME gen_regression_simple_driver_fortran
         COMMAND $<TARGET_FILE:gen_regression_simple_driver_fortran>)
-    set_target_properties(gen_regression_simple_driver_fortran PROPERTIES LINKER_LANGUAGE Fortran)
 endif()
 
 if(CMAKE_C_COMPILER_LOADED)
