norm: fix complex abs( inf + inf*i ) in CUDA and HIP

icl-utk-edu · Jul 7, 2024 · d47e9d6 · d47e9d6
1 parent 3e2934e
commit d47e9d6
Show file tree

Hide file tree

Showing 3 changed files with 11 additions and 11 deletions.
diff --git a/src/cuda/device_util.cuh b/src/cuda/device_util.cuh
@@ -158,7 +158,7 @@ inline float abs(cuFloatComplex x)
     // CUDA has a good implementation.
     return cuCabsf(x);
 #else
-    // For HIP, use our implementation that scales per LAPACK.
+    // For HIP, use our implementation that scales per LAPACK lapy2.
     float a = real( x );
     float b = imag( x );
     float z, w, t;
@@ -173,7 +173,7 @@ inline float abs(cuFloatComplex x)
         b = fabsf(b);
         w = max(a, b);
         z = min(a, b);
-        if (z == 0) {
+        if (z == 0 || isinf( w )) {
             t = w;
         }
         else {
@@ -194,7 +194,7 @@ inline double abs(cuDoubleComplex x)
     // CUDA has a good implementation.
     return cuCabs(x);
 #else
-    // For HIP, use our implementation that scales per LAPACK.
+    // For HIP, use our implementation that scales per LAPACK lapy2.
     double a = real( x );
     double b = imag( x );
     double z, w, t;
@@ -209,12 +209,12 @@ inline double abs(cuDoubleComplex x)
         b = fabs(b);
         w = max(a, b);
         z = min(a, b);
-        if (z == 0) {
+        if (z == 0 || isinf( w )) {
             t = w;
         }
         else {
             t = z/w;
-            t = 1.0 + t*t;
+            t = 1 + t*t;
             t = w * sqrt(t);
         }
         return t;

diff --git a/src/hip/device_util.hip.hh b/src/hip/device_util.hip.hh
@@ -158,7 +158,7 @@ inline float abs(rocblas_float_complex x)
     // CUDA has a good implementation.
     return hipCabsf(x);
 #else
-    // For HIP, use our implementation that scales per LAPACK.
+    // For HIP, use our implementation that scales per LAPACK lapy2.
     float a = real( x );
     float b = imag( x );
     float z, w, t;
@@ -173,7 +173,7 @@ inline float abs(rocblas_float_complex x)
         b = fabsf(b);
         w = max(a, b);
         z = min(a, b);
-        if (z == 0) {
+        if (z == 0 || isinf( w )) {
             t = w;
         }
         else {
@@ -194,7 +194,7 @@ inline double abs(rocblas_double_complex x)
     // CUDA has a good implementation.
     return hipCabs(x);
 #else
-    // For HIP, use our implementation that scales per LAPACK.
+    // For HIP, use our implementation that scales per LAPACK lapy2.
     double a = real( x );
     double b = imag( x );
     double z, w, t;
@@ -209,12 +209,12 @@ inline double abs(rocblas_double_complex x)
         b = fabs(b);
         w = max(a, b);
         z = min(a, b);
-        if (z == 0) {
+        if (z == 0 || isinf( w )) {
             t = w;
         }
         else {
             t = z/w;
-            t = 1.0 + t*t;
+            t = 1 + t*t;
             t = w * sqrt(t);
         }
         return t;

diff --git a/src/hip/device_util.hip.hh.dep b/src/hip/device_util.hip.hh.dep
@@ -1 +1 @@
-a403f5b1c0195b503252fa643c722495  src/cuda/device_util.cuh
+9aa8108d505dd3a8afee23324e45bb95  src/cuda/device_util.cuh
Original file line number	Diff line number	Diff line change
		@@ -1 +1 @@
		a403f5b1c0195b503252fa643c722495 src/cuda/device_util.cuh
		9aa8108d505dd3a8afee23324e45bb95 src/cuda/device_util.cuh