From: Daniel Arndt Date: Wed, 31 Oct 2018 18:41:09 +0000 (+0100) Subject: Change position of template argument in CUDAWrappers::kernel X-Git-Tag: v9.1.0-rc1~464^2~7 X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=4017e6d98b073a4d6005d64763ab165da4198e10;p=dealii.git 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 4059cd8a06..b375af82bd 100644 --- a/include/deal.II/lac/cuda_kernels.h +++ b/include/deal.II/lac/cuda_kernels.h @@ -26,6 +26,8 @@ # include +# include + DEAL_II_NAMESPACE_OPEN namespace LinearAlgebra @@ -57,9 +59,9 @@ namespace LinearAlgebra * * @ingroup CUDAWrappers */ + template struct Binop_Addition { - template __device__ static inline Number operation(const Number a, const Number b) { @@ -67,6 +69,18 @@ namespace LinearAlgebra } }; + template + struct Binop_Addition> + { + __device__ static inline std::complex + operation(const std::complex a, const std::complex b) + { + printf("This function is not implemented for std::complex!"); + assert(false); + return {}; + } + }; + /** @@ -74,9 +88,9 @@ namespace LinearAlgebra * * @ingroup CUDAWrappers */ + template struct Binop_Subtraction { - template __device__ static inline Number operation(const Number a, const Number b) { @@ -91,9 +105,9 @@ namespace LinearAlgebra * * @ingroup CUDAWrappers */ - template + template 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 3a22883453..7a8185f302 100644 --- 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 + template 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::operation(v1[idx], v2[idx]); } } diff --git a/source/lac/cuda_kernels.cu b/source/lac/cuda_kernels.cu index 1ec88bd11d..49230c517c 100644 --- a/source/lac/cuda_kernels.cu +++ b/source/lac/cuda_kernels.cu @@ -31,11 +31,11 @@ namespace LinearAlgebra vec_scale(float *, const float a, const size_type); template __global__ void vector_bin_op(float * v1, - float * v2, + const float * v2, const size_type N); template __global__ void vector_bin_op(float * v1, - float * v2, + const float * v2, const size_type N); template struct ElemSum; template struct L1Norm; @@ -126,11 +126,11 @@ namespace LinearAlgebra vec_scale(double *, const double a, const size_type); template __global__ void vector_bin_op(double * v1, - double * v2, + const double * v2, const size_type N); template __global__ void vector_bin_op(double * v1, - double * v2, + const double * v2, const size_type N); template struct ElemSum; template struct L1Norm;