Patch addAlpha definitions to work around buggy HIP complex operator …

…overloads
eth-cscs · Sep 25, 2024 · 0ece43f · 0ece43f
1 parent 0208a04
commit 0ece43f
Showing 1 changed file with 12 additions and 0 deletions.
diff --git a/src/lapack/gpu/add.cu b/src/lapack/gpu/add.cu
@@ -32,6 +32,18 @@ __device__ inline void addAlpha(const T& alpha, const T& a, T& b) {
   b = b + alpha * a;
 }
 
+template <>
+__device__ inline void addAlpha<hipFloatComplex>(const hipFloatComplex& alpha, const hipFloatComplex& a,
+                                                 hipFloatComplex& b) {
+  b = b + hipCmulf(alpha, a);
+}
+
+template <>
+__device__ inline void addAlpha<hipDoubleComplex>(const hipDoubleComplex& alpha,
+                                                  const hipDoubleComplex& a, hipDoubleComplex& b) {
+  b = b + hipCmul(alpha, a);
+}
+
 template <class T>
 __device__ inline void sum(const T& /*alpha*/, const T& a, T& b) {
   b = b + a;