From af35d1bfeb36511a5c2bbda4647370505a4d6b65 Mon Sep 17 00:00:00 2001
From: Kaushik Kulkarni <kaushikcfd@gmail.com>
Date: Sun, 7 Aug 2022 12:19:10 -0500
Subject: [PATCH 1/2] use constexpr for dynamic instantiation of constants

CUDA doesn't support static keywords for device kernels.
---
 pycuda/cuda/pycuda-complex-impl.hpp | 13 ++++++-------
 1 file changed, 6 insertions(+), 7 deletions(-)

diff --git a/pycuda/cuda/pycuda-complex-impl.hpp b/pycuda/cuda/pycuda-complex-impl.hpp
index d2e67605..68e6f01a 100644
--- a/pycuda/cuda/pycuda-complex-impl.hpp
+++ b/pycuda/cuda/pycuda-complex-impl.hpp
@@ -173,9 +173,12 @@ __device__ complex<float>  exp(const complex<float>& z)
 __device__ complex<double> exp(const complex<double>& z)
 { return expT(z); }
 
-#if 0
 //----------------------------------------------------------------------
 // log10
+
+__device__ constexpr float log10_invf() { return 1.f / ::log(10.f); }
+__device__ constexpr double log10_inv() { return 1.0 / ::log(10.); }
+
 template <class _Tp>
 static __device__ complex<_Tp> log10T(const complex<_Tp>& z, const _Tp& ln10_inv) {
   complex<_Tp> r;
@@ -185,15 +188,11 @@ static __device__ complex<_Tp> log10T(const complex<_Tp>& z, const _Tp& ln10_inv
   return r;
 }
 
-static const float LN10_INVF = 1.f / ::log(10.f);
 __device__ complex<float> log10(const complex<float>& z)
-{ return log10T(z, LN10_INVF); }
+{ return log10T(z, log10_invf()); }
 
-static const double LN10_INV = 1. / ::log10(10.);
 __device__ complex<double> log10(const complex<double>& z)
-{ return log10T(z, LN10_INV); }
-
-#endif
+{ return log10T(z, log10_inv()); }
 
 //----------------------------------------------------------------------
 // log
-- 
GitLab


From 09fd4c3237bbf5d9958ef45c1cde8064a90bd0e3 Mon Sep 17 00:00:00 2001
From: Kaushik Kulkarni <kaushikcfd@gmail.com>
Date: Sun, 7 Aug 2022 12:24:32 -0500
Subject: [PATCH 2/2] test log10

---
 test/test_gpuarray.py | 11 +++++++++++
 1 file changed, 11 insertions(+)

diff --git a/test/test_gpuarray.py b/test/test_gpuarray.py
index cce2de92..d4e03f7d 100644
--- a/test/test_gpuarray.py
+++ b/test/test_gpuarray.py
@@ -1380,6 +1380,17 @@ class TestGPUArray:
         a[...] = 1729
         np.testing.assert_allclose(a.get(), 1729)
 
+    @pytest.mark.parametrize("dtype,rtol", [(np.complex64, 1e-6),
+                                            (np.complex128, 1e-14)])
+    def test_log10(self, dtype, rtol):
+        from pycuda import cumath
+
+        rng = np.random.default_rng(seed=0)
+        x_np = rng.random((10, 4)) + dtype(1j)*rng.random((10, 4))
+        x_cu = gpuarray.to_gpu(x_np)
+        np.testing.assert_allclose(cumath.log10(x_cu).get(), np.log10(x_np),
+                                   rtol=rtol)
+
 
 if __name__ == "__main__":
     # make sure that import failures get reported, instead of skipping the tests.
-- 
GitLab