220 mgis functions add support for nvcc

CMakeLists.txt

@@ -39,7 +39,19 @@ option(enable-mgis-function-precompiled-operations
 option(enable-exceptions "use exceptions to report contract violation and error reporting" OFF)
 
 # support for STL parallel algorithms
-option(enable-parallel-stl-algorithms "enable STL parallel algorithms" ON)
+option(enable-parallel-stl-algorithms "enable STL parallel algorithms" OFF)
+
+# support for CUDA algorithms
+option(enable-cuda-algorithms "enable CUDA parallel algorithms" OFF)
+
+if(enable-cuda-algorithms)
+  enable_language(CUDA)
+  if(CMAKE_CUDA_COMPILER_ID STREQUAL "NVIDIA")
+    set(MGIS_CUDA_FLAGS "--expt-relaxed-constexpr" "--extended-lambda" "--diag-suppress=20012")
+  else()
+    set(MGIS_CUDA_FLAGS)
+  endif(CMAKE_CUDA_COMPILER_ID STREQUAL "NVIDIA")
+endif(enable-cuda-algorithms)
 
 # C-bindings (placed before compiler detection)
 option(enable-c-bindings "enable c bindings support")

INSTALL-cmake.md

@@ -66,6 +66,8 @@ Options
 - `enable-exceptions`: use exceptions to report contract violation and error reporting.
    By default, contract violation leads to abort the program.
 - `enable-parallel-stl-algorithms`: by default, STL algorithms are used if available
+- `enable-cuda-algorithms`: add support for algorithms expressed in
+  `CUDA` programming model.
 - `enable-hdf5-support`: enable HDF5 support for save/restore
   operations. Note that if this support is not explicitely requested,
   `MGIS` still tries by default to support `HDF5` but configuration will

docs/web/release-notes-3.2.md

@@ -24,6 +24,11 @@ This version is meant to be used with `TFEL` Version 5.2.
 
 ## New `cmake` options
 
+### Support for `CUDA` in `MGIS/Function`
+
+Support for `CUDA` in `MGIS/Function` is enable by
+`enable-cuda-algorithms`.
+
 ### HDF5 support
 
 The following options control the support of the `HDF5` library:
@@ -181,6 +186,10 @@ const auto e2 = f.get<0, tfel::math::stensor<2, real>>(1);
 
 # Issues fixed
 
+## Issue 220: [mgis-functions] Add support for `nvcc`
+
+For more details, see <https://github.com/thelfer/MFrontGenericInterfaceSupport/issues/220>
+
 ## Issue 210: [documentation] Deployment of the doxygen documentation
 
 For more details, see <https://github.com/thelfer/MFrontGenericInterfaceSupport/issues/210>

include/CMakeLists.txt

@@ -100,6 +100,10 @@ if(enable-mgis-function)
   mgis_header(MGIS/Function CoalescedMemoryAccessFunctionViewBase.ixx)
   mgis_header(MGIS/Function StridedCoalescedMemoryAccessFunctionViewBase.hxx)
   mgis_header(MGIS/Function StridedCoalescedMemoryAccessFunctionViewBase.ixx)
+  if(enable-cuda-algorithms)
+    mgis_header(MGIS/Function/CUDA Algorithms.hxx)
+    mgis_header(MGIS/Function/CUDA Algorithms.ixx)
+  endif(enable-cuda-algorithms)
   if(MGIS_HAVE_TFEL)
     mgis_header(MGIS/Function/TFEL TensorConcept.hxx)
     mgis_header(MGIS/Function/TFEL Tensors.hxx)

include/MGIS/Config-c.h

@@ -143,4 +143,23 @@ typedef size_t mgis_size_type;
 #define MGIS_TEMPORARY_VARIABLE(X) \
   MGIS_PP_JOIN(MGIS_PP_JOIN(MGIS_PP_JOIN(mgis_temporary_, X), _), __LINE__)
 
+#if (defined __CUDACC__) || (defined __HIPCC__)
+#ifndef MGIS_DEVICE
+#define MGIS_DEVICE __device__
+#endif /* MGIS_DEVICE */
+#ifndef MGIS_HOST
+#define MGIS_HOST __host__
+#endif /* MGIS_HOST */
+#endif /* (defined __CUDACC__) || (defined __HIPCC__) */
+
+#ifndef MGIS_DEVICE
+#define MGIS_DEVICE
+#endif /* MGIS_DEVICE */
+
+#ifndef MGIS_HOST
+#define MGIS_HOST
+#endif /* MGIS_HOST */
+
+#define MGIS_HOST_DEVICE MGIS_DEVICE MGIS_HOST
+
 #endif /* LIB_MGIS_CONFIG_C_H */

include/MGIS/Function/BinaryOperation.ixx

@@ -23,9 +23,9 @@ namespace mgis::function {
   requires(BinaryOperationModifierRequirements<CallableType,
                                                FirstEvaluatorType,
                                                SecondEvaluatorType>)  //
-      constexpr BinaryOperationModifier<CallableType,
-                                        FirstEvaluatorType,
-                                        SecondEvaluatorType>::
+      MGIS_HOST_DEVICE constexpr BinaryOperationModifier<CallableType,
+                                                         FirstEvaluatorType,
+                                                         SecondEvaluatorType>::
           BinaryOperationModifier(const CallableType& c,
                                   const FirstEvaluatorType& e1,
                                   const SecondEvaluatorType& e2)
@@ -40,6 +40,7 @@ namespace mgis::function {
   requires(BinaryOperationModifierRequirements<CallableType,
                                                FirstEvaluatorType,
                                                SecondEvaluatorType>)  //
+      MGIS_HOST_DEVICE
       constexpr auto BinaryOperationModifier<CallableType,
                                              FirstEvaluatorType,
                                              SecondEvaluatorType>::
@@ -54,6 +55,7 @@ namespace mgis::function {
   requires(BinaryOperationModifier2Requirements<CallableType,
                                                 FirstEvaluatorType,
                                                 SecondEvaluatorType>)  //
+      MGIS_HOST_DEVICE
       constexpr auto BinaryOperationModifier2<CallableType,
                                               FirstEvaluatorType,
                                               SecondEvaluatorType>::
@@ -66,7 +68,7 @@ namespace mgis::function {
   template <typename CallableType,
             EvaluatorConcept FirstEvaluatorType,
             EvaluatorConcept SecondEvaluatorType>
-  constexpr mgis::size_type getNumberOfComponents(
+  MGIS_HOST_DEVICE constexpr mgis::size_type getNumberOfComponents(
       const BinaryOperationModifier<CallableType,
                                     FirstEvaluatorType,
                                     SecondEvaluatorType>& e) {
@@ -76,7 +78,7 @@ namespace mgis::function {
   template <typename CallableType,
             EvaluatorConcept FirstEvaluatorType,
             EvaluatorConcept SecondEvaluatorType>
-  constexpr mgis::size_type getNumberOfComponents(
+  MGIS_HOST_DEVICE constexpr mgis::size_type getNumberOfComponents(
       const BinaryOperationModifier2<CallableType,
                                      FirstEvaluatorType,
                                      SecondEvaluatorType>& e) {
@@ -87,6 +89,7 @@ namespace mgis::function {
 
     template <typename CallableType, EvaluatorConcept SecondEvaluatorType>
     requires(std::is_copy_constructible_v<CallableType>)  //
+        MGIS_HOST_DEVICE
         constexpr BinaryOperatorCurrying<CallableType, SecondEvaluatorType>::
             BinaryOperatorCurrying(const CallableType& c,
                                    const SecondEvaluatorType& e)
@@ -95,7 +98,7 @@ namespace mgis::function {
     template <typename CallableType, EvaluatorConcept SecondEvaluatorType>
     requires(std::is_copy_constructible_v<CallableType>)  //
         template <typename FirstEvaluatorType>
-        constexpr auto BinaryOperatorCurrying<
+        MGIS_HOST_DEVICE constexpr auto BinaryOperatorCurrying<
             CallableType,
             SecondEvaluatorType>::operator()(FirstEvaluatorType&& e1) const
         requires((EvaluatorConcept<std::decay_t<FirstEvaluatorType>>)&&(
@@ -109,32 +112,36 @@ namespace mgis::function {
 
     template <typename CallableType, EvaluatorConcept SecondEvaluatorType>
     requires(std::is_trivially_default_constructible_v<CallableType>)  //
+        MGIS_HOST_DEVICE
         constexpr BinaryOperatorCurrying2<CallableType, SecondEvaluatorType>::
             BinaryOperatorCurrying2(const SecondEvaluatorType& e)
         : e2(e) {}
 
     template <typename CallableType, EvaluatorConcept SecondEvaluatorType>
-    requires(std::is_trivially_default_constructible_v<CallableType>) template <
-        typename FirstEvaluatorType>
-    constexpr auto BinaryOperatorCurrying2<CallableType, SecondEvaluatorType>::
-    operator()(FirstEvaluatorType&& e1) const requires(
-        (EvaluatorConcept<std::decay_t<FirstEvaluatorType>>)&&  //
-        (std::invocable<CallableType,
-                        evaluator_result<std::decay_t<FirstEvaluatorType>>,
-                        evaluator_result<SecondEvaluatorType>>)) {
+    requires(std::is_trivially_default_constructible_v<CallableType>)  //
+        template <typename FirstEvaluatorType>
+        MGIS_HOST_DEVICE constexpr auto BinaryOperatorCurrying2<
+            CallableType,
+            SecondEvaluatorType>::operator()(FirstEvaluatorType&& e1) const
+        requires(
+            (EvaluatorConcept<std::decay_t<FirstEvaluatorType>>)&&  //
+            (std::invocable<CallableType,
+                            evaluator_result<std::decay_t<FirstEvaluatorType>>,
+                            evaluator_result<SecondEvaluatorType>>)) {
       return BinaryOperationModifier2<
           CallableType, std::decay_t<FirstEvaluatorType>, SecondEvaluatorType>(
           std::forward<FirstEvaluatorType>(e1), this->e2);
     }  // end of operator()
 
     template <typename CallableType>
-    constexpr binary_operation_modifier<
+    MGIS_HOST_DEVICE constexpr binary_operation_modifier<
         CallableType>::binary_operation_modifier(const CallableType& c)
         : modifier(c) {}  // end of binary_operation_modifier
 
     template <typename CallableType>
     template <typename SecondEvaluatorType>
-    constexpr auto binary_operation_modifier<CallableType>::operator()(
+    MGIS_HOST_DEVICE constexpr auto
+    binary_operation_modifier<CallableType>::operator()(
         SecondEvaluatorType&& e2) const
         requires(EvaluatorConcept<std::decay_t<SecondEvaluatorType>>) {
       return BinaryOperatorCurrying<CallableType,
@@ -144,7 +151,8 @@ namespace mgis::function {
 
     template <typename CallableType>
     template <typename FirstEvaluatorType, typename SecondEvaluatorType>
-    constexpr auto binary_operation_modifier<CallableType>::operator()(
+    MGIS_HOST_DEVICE constexpr auto
+    binary_operation_modifier<CallableType>::operator()(
         FirstEvaluatorType&& e1, SecondEvaluatorType&& e2) const
         requires((EvaluatorConcept<std::decay_t<FirstEvaluatorType>>)&&   //
                  (EvaluatorConcept<std::decay_t<SecondEvaluatorType>>)&&  //
@@ -161,7 +169,8 @@ namespace mgis::function {
 
     template <typename CallableType>
     template <typename SecondEvaluatorType>
-    constexpr auto binary_operation_modifier2_impl<CallableType>::operator()(
+    MGIS_HOST_DEVICE constexpr auto
+    binary_operation_modifier2_impl<CallableType>::operator()(
         SecondEvaluatorType&& e2) const
         requires(EvaluatorConcept<std::decay_t<SecondEvaluatorType>>) {
       return BinaryOperatorCurrying2<CallableType,
@@ -171,7 +180,8 @@ namespace mgis::function {
 
     template <typename CallableType>
     template <typename FirstEvaluatorType, typename SecondEvaluatorType>
-    constexpr auto binary_operation_modifier2_impl<CallableType>::operator()(
+    MGIS_HOST_DEVICE constexpr auto
+    binary_operation_modifier2_impl<CallableType>::operator()(
         FirstEvaluatorType&& e1, SecondEvaluatorType&& e2) const
         requires((EvaluatorConcept<std::decay_t<FirstEvaluatorType>>)&&   //
                  (EvaluatorConcept<std::decay_t<SecondEvaluatorType>>)&&  //
@@ -187,15 +197,15 @@ namespace mgis::function {
     }  // end of operator()
 
     template <typename CallableType>
-    constexpr auto binary_operation_modifier2(CallableType) {
+    MGIS_HOST_DEVICE constexpr auto binary_operation_modifier2(CallableType) {
       return binary_operation_modifier2_impl<CallableType>{};
     }  // end of binary_operation_modifier2
 
   }  // namespace internals
 
   template <typename CallableType, typename SecondEvaluatorType>
-  constexpr auto binary_operation(CallableType&& c,
-                                  SecondEvaluatorType&& e2)  //
+  MGIS_HOST_DEVICE constexpr auto binary_operation(CallableType&& c,
+                                                   SecondEvaluatorType&& e2)  //
       requires(EvaluatorConcept<std::decay_t<SecondEvaluatorType>>) {
     auto modifier =
         internals::binary_operation_modifier<std::decay_t<CallableType>>(

include/MGIS/Function/CUDA/Algorithms.hxx

@@ -0,0 +1,56 @@
+/*!
+ * \file   MGIS/Function/CUDA/Algorithms.hxx
+ * \brief
+ * \author Thomas Helfer
+ * \date   28/06/2026
+ * \copyright (C) Copyright Thomas Helfer 2018.
+ * Use, modification and distribution are subject
+ * to one of the following licences:
+ * - GNU Lesser General Public License (LGPL), Version 3.0. (See accompanying
+ *   file LGPL-3.0.txt)
+ * - CECILL-C,  Version 1.0 (See accompanying files
+ *   CeCILL-C_V1-en.txt and CeCILL-C_V1-fr.txt).
+ */
+
+#ifndef LIB_MGIS_FUNCTION_CUDA_ALGORITHMS_HXX
+#define LIB_MGIS_FUNCTION_CUDA_ALGORITHMS_HXX
+
+#include "MGIS/Config.hxx"
+#include "MGIS/Context.hxx"
+#include "MGIS/Function/SpaceConcept.hxx"
+#include "MGIS/Function/EvaluatorConcept.hxx"
+#include "MGIS/Function/FunctionConcept.hxx"
+
+namespace mgis::function {
+
+  /*!
+   * \brief execution configuration used to launch CUDA's kernels
+   */
+  struct CUDAExecutionConfiguration {
+    //! \brief number of threads per block
+    int number_of_threads_per_block = 32;
+  };
+
+  /*!
+   * \brief assign the evaluator to a function
+   * \param[in] ctx: execution context
+   * \param[in] c: CUDA execution configuration
+   * \param[in] lhs: left hand side
+   * \param[in] e: right hand side
+   */
+  template <typename FunctionType, EvaluatorConcept EvaluatorType>
+  [[nodiscard]] bool assign(AbstractErrorHandler&,
+                            const CUDAExecutionConfiguration&,
+                            FunctionType&,
+                            const EvaluatorType)  //
+      requires(
+          ((LinearElementSpaceConcept<evaluator_space<EvaluatorType>>) ||
+           (LinearQuadratureSpaceConcept<evaluator_space<EvaluatorType>>)) &&
+          std::same_as<function_space<FunctionType>,
+                       evaluator_space<EvaluatorType>>);
+
+}  // end of namespace mgis::function
+
+#include "MGIS/Function/CUDA/Algorithms.ixx"
+
+#endif /* LIB_MGIS_FUNCTION_CUDA_ALGORITHMS_HXX */