Change position of template argument in CUDAWrappers::kernel

diff --git a/include/deal.II/lac/cuda_kernels.h b/include/deal.II/lac/cuda_kernels.h
index 4059cd8a06fd79211fd3fa5a0a7e3ac63543c9b8..b375af82bde3e2c9d54c702b5f2552f72280cc07 100644 (file)
--- a/include/deal.II/lac/cuda_kernels.h
+++ b/include/deal.II/lac/cuda_kernels.h
@@ -26,6 +26,8 @@
 
 #  include <deal.II/lac/cuda_atomic.h>
 
+#  include <assert.h>
+
 DEAL_II_NAMESPACE_OPEN
 
 namespace LinearAlgebra
@@ -57,9 +59,9 @@ namespace LinearAlgebra
        *
        * @ingroup CUDAWrappers
        */
+      template <typename Number>
       struct Binop_Addition
       {
-        template <typename Number>
         __device__ static inline Number
         operation(const Number a, const Number b)
         {
@@ -67,6 +69,18 @@ namespace LinearAlgebra
         }
       };
 
+      template <typename Number>
+      struct Binop_Addition<std::complex<Number>>
+      {
+        __device__ static inline std::complex<Number>
+        operation(const std::complex<Number> a, const std::complex<Number> b)
+        {
+          printf("This function is not implemented for std::complex<Number>!");
+          assert(false);
+          return {};
+        }
+      };
+
 
 
       /**
@@ -74,9 +88,9 @@ namespace LinearAlgebra
        *
        * @ingroup CUDAWrappers
        */
+      template <typename Number>
       struct Binop_Subtraction
       {
-        template <typename Number>
         __device__ static inline Number
         operation(const Number a, const Number b)
         {
@@ -91,9 +105,9 @@ namespace LinearAlgebra
        *
        * @ingroup CUDAWrappers
        */
-      template <typename Number, typename Binop>
+      template <typename Number, template <typename> class Binop>
       __global__ void
-      vector_bin_op(Number *v1, Number *v2, const size_type N);
+      vector_bin_op(Number *v1, const Number *v2, const size_type N);
 
 
 

diff --git a/include/deal.II/lac/cuda_kernels.templates.h b/include/deal.II/lac/cuda_kernels.templates.h
index 3a228834535c04dc7efb0aee27b73ef788049f40..7a8185f302e22eba946dfda5e9e738102fc786d5 100644 (file)
--- a/include/deal.II/lac/cuda_kernels.templates.h
+++ b/include/deal.II/lac/cuda_kernels.templates.h
@@ -44,9 +44,9 @@ namespace LinearAlgebra
 
 
 
-      template <typename Number, typename Binop>
+      template <typename Number, template <typename> class Binop>
       __global__ void
-      vector_bin_op(Number *v1, Number *v2, const size_type N)
+      vector_bin_op(Number *v1, const Number *v2, const size_type N)
       {
         const size_type idx_base =
           threadIdx.x + blockIdx.x * (blockDim.x * chunk_size);
@@ -54,7 +54,7 @@ namespace LinearAlgebra
           {
             const size_type idx = idx_base + i * block_size;
             if (idx < N)
-              v1[idx] = Binop::operation(v1[idx], v2[idx]);
+              v1[idx] = Binop<Number>::operation(v1[idx], v2[idx]);
           }
       }
 

diff --git a/source/lac/cuda_kernels.cu b/source/lac/cuda_kernels.cu
index 1ec88bd11dc8ecbe4cebeb51324016cb01db63b5..49230c517cc3339e239c302e79c2be86233c1b72 100644 (file)
--- a/source/lac/cuda_kernels.cu
+++ b/source/lac/cuda_kernels.cu
@@ -31,11 +31,11 @@ namespace LinearAlgebra
       vec_scale<float>(float *, const float a, const size_type);
       template __global__ void
       vector_bin_op<float, Binop_Addition>(float *         v1,
-                                           float *         v2,
+                                           const float *   v2,
                                            const size_type N);
       template __global__ void
       vector_bin_op<float, Binop_Subtraction>(float *         v1,
-                                              float *         v2,
+                                              const float *   v2,
                                               const size_type N);
       template struct ElemSum<float>;
       template struct L1Norm<float>;
@@ -126,11 +126,11 @@ namespace LinearAlgebra
       vec_scale<double>(double *, const double a, const size_type);
       template __global__ void
       vector_bin_op<double, Binop_Addition>(double *        v1,
-                                            double *        v2,
+                                            const double *  v2,
                                             const size_type N);
       template __global__ void
       vector_bin_op<double, Binop_Subtraction>(double *        v1,
-                                               double *        v2,
+                                               const double *  v2,
                                                const size_type N);
       template struct ElemSum<double>;
       template struct L1Norm<double>;