From b740a667be720017820d0f9f29512640458a83ca Mon Sep 17 00:00:00 2001 From: Mark Gates Date: Sun, 7 Jul 2024 00:00:12 -0400 Subject: [PATCH] fix complex abs( inf + inf*i ) in CUDA and HIP --- src/cuda/device_util.cuh | 6 +++--- src/hip/device_util.hip.hh | 6 +++--- src/hip/device_util.hip.hh.dep | 2 +- 3 files changed, 7 insertions(+), 7 deletions(-) diff --git a/src/cuda/device_util.cuh b/src/cuda/device_util.cuh index 2170afc2..1d8bc019 100644 --- a/src/cuda/device_util.cuh +++ b/src/cuda/device_util.cuh @@ -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,7 +209,7 @@ 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 { diff --git a/src/hip/device_util.hip.hh b/src/hip/device_util.hip.hh index a2f89aa7..e8c3018f 100644 --- a/src/hip/device_util.hip.hh +++ b/src/hip/device_util.hip.hh @@ -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,7 +209,7 @@ 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 { diff --git a/src/hip/device_util.hip.hh.dep b/src/hip/device_util.hip.hh.dep index c914bb9e..bec5c324 100644 --- a/src/hip/device_util.hip.hh.dep +++ b/src/hip/device_util.hip.hh.dep @@ -1 +1 @@ -a403f5b1c0195b503252fa643c722495 src/cuda/device_util.cuh +3a2509f9538de94f6212c4e678b236a6 src/cuda/device_util.cuh