Patch addAlpha definitions to work around buggy HIP complex operator …

…overloads
eth-cscs · Sep 26, 2024 · 52f7ee2 · 52f7ee2
1 parent 0208a04
commit 52f7ee2
Showing 1 changed file with 14 additions and 0 deletions.
diff --git a/src/lapack/gpu/add.cu b/src/lapack/gpu/add.cu
@@ -32,6 +32,20 @@ __device__ inline void addAlpha(const T& alpha, const T& a, T& b) {
   b = b + alpha * a;
 }
 
+#ifdef DLAF_WITH_HIP
+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);
+}
+#endif
+
 template <class T>
 __device__ inline void sum(const T& /*alpha*/, const T& a, T& b) {
   b = b + a;