EXSWHTEC-286 - Implement tests for log device math functions #230

Change-Id: I642a9d865fcc30d7b303b0d4dd05fcd723a59015 [ROCm/hip-tests commit: 95a75cf00f]
2024-01-22 23:46:23 +05:30
@@ -173,7 +173,7 @@ if (WIN32)
 endif()

 if(HIP_PLATFORM STREQUAL "amd")
-    add_compile_options(-Wall -Wextra -Wvla -Werror -Wno-deprecated -Wno-option-ignored)
+    add_compile_options(-Wall -Wextra -Wvla -Werror -Wno-deprecated -Wno-option-ignored -Wno-unused-parameter -Wunused-variable)
 endif()

 cmake_policy(PUSH)
@@ -131,6 +131,9 @@
        "=== Patch which removes the typetraits implementation from std namespace in hiprtc is reverted ===",
        "Unit_hiprtc_stdheaders",
        "Unit_hipGraphAddMemcpyNode_Negative_Parameters",
+        "=== Below 2 tests are disable due to defect EXSWHTEC-369 ===",
+        "Unit_Device_ilogbf_Accuracy_Positive",
+        "Unit_Device_ilogb_Accuracy_Positive",
        "Unit_hipMemAddressFree_negative",
        "Unit_hipMemAddressReserve_AlignmentTest",
        "Unit_hipMemAddressReserve_Negative",
@@ -222,6 +222,9 @@
        "NOTE: The following test is disabled due to defect - EXSWHTEC-244",
        "Unit_hipExtLaunchMultiKernelMultiDevice_Negative_Parameters",
        "Unit_hipMemAddressFree_negative",
+        "=== Below 2 tests are disable due to defect EXSWHTEC-369 ===",
+        "Unit_Device_ilogbf_Accuracy_Positive",
+        "Unit_Device_ilogb_Accuracy_Positive",
        "Unit_hipMemAddressReserve_AlignmentTest",
        "Unit_hipGraphAddMemcpyNode_Negative_Parameters",
        "Unit_hipMemCreate_ChkWithKerLaunch",
@@ -27,6 +27,7 @@ set(TEST_SRC
  integer_intrinsics.cc
  root_funcs.cc
  pow_funcs.cc
+  log_funcs.cc
 )

 if(HIP_PLATFORM MATCHES "nvidia")
@@ -91,3 +92,7 @@ add_test(NAME Unit_Device_pow_Negative
         COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py
         ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH}
         math_pow_negative_kernels.cc 76)
+add_test(NAME Unit_Device_log_Negative
+         COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py
+         ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH}
+         math_log_negative_kernels.cc 24)
@@ -0,0 +1,260 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "unary_common.hh"
+#include "math_log_negative_kernels_rtc.hh"
+
+/**
+ * @addtogroup LogMathFuncs LogMathFuncs
+ * @{
+ * @ingroup MathTest
+ */
+
+/********** Unary Functions **********/
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Tests the numerical accuracy of `logf(x)` for all possible inputs and `log(x)` against a
+ * table of difficult values, followed by a large number of randomly generated values. The results
+ * are compared against reference function `T std::log(T)`. The maximum ulp error is 1.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/log_funcs.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+MATH_UNARY_WITHIN_ULP_STL_REF_TEST_DEF(log, 1, 1)
+
+/**
+ * Test Description
+ * ------------------------
+ *    - RTCs kernels that pass argument of invalid type for logf and log.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/log_funcs.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device_log_logf_Negative_RTC") { NegativeTestRTCWrapper<4>(kLog); }
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Tests the numerical accuracy of `log2f(x)` for all possible inputs and `log2(x)` against a
+ * table of difficult values, followed by a large number of randomly generated values. The results
+ * are compared against reference function `T std::log2(T)`. The maximum ulp error is 1.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/log_funcs.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+MATH_UNARY_WITHIN_ULP_STL_REF_TEST_DEF(log2, 1, 1)
+
+/**
+ * Test Description
+ * ------------------------
+ *    - RTCs kernels that pass argument of invalid type for log2f and log2.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/log_funcs.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device_log2_log2f_Negative_RTC") { NegativeTestRTCWrapper<4>(kLog2); }
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Tests the numerical accuracy of `log10f(x)` for all possible inputs and `log10(x)` against a
+ * table of difficult values, followed by a large number of randomly generated values. The results
+ * are compared against reference function `T std::log10(T)`. The maximum ulp error for single
+ * precision is 2 and for double precision is 1.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/log_funcs.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+MATH_UNARY_WITHIN_ULP_STL_REF_TEST_DEF(log10, 2, 1)
+
+/**
+ * Test Description
+ * ------------------------
+ *    - RTCs kernels that pass argument of invalid type for log10f and log10.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/log_funcs.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device_log10_log10f_Negative_RTC") { NegativeTestRTCWrapper<4>(kLog10); }
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Tests the numerical accuracy of `log1pf(x)` for all possible inputs and `log1p(x)` against a
+ * table of difficult values, followed by a large number of randomly generated values. The results
+ * are compared against reference function `T std::log1p(T)`. The maximum ulp error is 1.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/log_funcs.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+MATH_UNARY_WITHIN_ULP_STL_REF_TEST_DEF(log1p, 1, 1)
+
+/**
+ * Test Description
+ * ------------------------
+ *    - RTCs kernels that pass argument of invalid type for log1pf and log1p.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/log_funcs.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device_log1p_log1pf_Negative_RTC") { NegativeTestRTCWrapper<4>(kLog1p); }
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Tests the numerical accuracy of `logb(x)` for all possible inputs and `logb(x)` against a
+ * table of difficult values, followed by a large number of randomly generated values. The results
+ * are compared against reference function `T std::logb(T)`. The maximum ulp error is 0.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/log_funcs.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+MATH_UNARY_WITHIN_ULP_STL_REF_TEST_DEF(logb, 0, 0)
+
+/**
+ * Test Description
+ * ------------------------
+ *    - RTCs kernels that pass argument of invalid type for logbf and logb.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/log_funcs.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device_logb_logbf_Negative_RTC") { NegativeTestRTCWrapper<4>(kLogb); }
+
+
+template <typename T>
+__global__ void ilogb_kernel(int* const ys, const size_t num_xs, T* const xs) {
+  const auto tid = cg::this_grid().thread_rank();
+  const auto stride = cg::this_grid().size();
+
+  for (auto i = tid; i < num_xs; i += stride) {
+    if constexpr (std::is_same_v<float, T>) {
+      ys[i] = ilogbf(xs[i]);
+    } else if constexpr (std::is_same_v<double, T>) {
+      ys[i] = ilogb(xs[i]);
+    }
+  }
+}
+
+template <typename T> int ilogb_ref(T arg) {
+  if (arg == 0) {
+    return std::numeric_limits<int>::min();
+  } else if (std::isnan(arg)) {
+    return std::numeric_limits<int>::min();
+  } else if (std::isinf(arg)) {
+    return std::numeric_limits<int>::max();
+  } else {
+    return std::ilogb(arg);
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Tests the numerical accuracy of `ilogbf(x)` for all possible inputs. The results are
+ * compared against reference function `int std::ilogb(double)`. The maximum ulp error is 0.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/log_funcs.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device_ilogbf_Accuracy_Positive") {
+  UnarySinglePrecisionTest(ilogb_kernel<float>, ilogb_ref<double>,
+                           EqValidatorBuilderFactory<int>());
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Tests the numerical accuracy of `ilogb(x)` against a table of difficult values,
+ * followed by a large number of randomly generated values. The results are
+ * compared against reference function `int std::ilogb(long double)`. The maximum ulp error is 0.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/log_funcs.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device_ilogb_Accuracy_Positive") {
+  UnaryDoublePrecisionTest(ilogb_kernel<double>, ilogb_ref<long double>,
+                           EqValidatorBuilderFactory<int>());
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - RTCs kernels that pass argument of invalid type for ilogbf and ilogb.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/log_funcs.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device_ilogb_ilogbf_Negative_RTC") { NegativeTestRTCWrapper<4>(kIlogb); }
@@ -0,0 +1,39 @@
+/*
+Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip_test_common.hh>
+
+class Dummy {
+ public:
+  __device__ Dummy() {}
+  __device__ ~Dummy() {}
+};
+
+#define NEGATIVE_KERNELS_SHELL(func_name)                                                          \
+  __global__ void func_name##_kernel_v1(double* x) { double result = func_name(x); }               \
+  __global__ void func_name##_kernel_v2(Dummy x) { double result = func_name(x); }                 \
+  __global__ void func_name##f_kernel_v1(float* x) { float result = func_name##f(x); }             \
+  __global__ void func_name##f_kernel_v2(Dummy x) { float result = func_name##f(x); }
+
+NEGATIVE_KERNELS_SHELL(log)
+NEGATIVE_KERNELS_SHELL(log2)
+NEGATIVE_KERNELS_SHELL(log10)
+NEGATIVE_KERNELS_SHELL(log1p)
+NEGATIVE_KERNELS_SHELL(logb)
+NEGATIVE_KERNELS_SHELL(ilogb)
@@ -0,0 +1,96 @@
+/*
+Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#pragma once
+
+/*
+Negative kernels used for the math log negative Test Cases that are using RTC.
+*/
+
+static constexpr auto kLog{R"(
+  class Dummy {
+   public:
+    __device__ Dummy() {}
+    __device__ ~Dummy() {}
+  };
+  __global__ void log_kernel_v1(double* x) { double result = log(x); }
+  __global__ void log_kernel_v2(Dummy x) { double result = log(x); }
+  __global__ void logf_kernel_v1(float* x) { float result = logf(x); }
+  __global__ void logf_kernel_v2(Dummy x) { float result = logf(x); }
+)"};
+
+static constexpr auto kLog2{R"(
+  class Dummy {
+   public:
+    __device__ Dummy() {}
+    __device__ ~Dummy() {}
+  };
+  __global__ void log2_kernel_v1(double* x) { double result = log2(x); }
+  __global__ void log2_kernel_v2(Dummy x) { double result = log2(x); }
+  __global__ void log2f_kernel_v1(float* x) { float result = log2f(x); }
+  __global__ void log2f_kernel_v2(Dummy x) { float result = log2f(x); }
+)"};
+
+static constexpr auto kLog10{R"(
+  class Dummy {
+   public:
+    __device__ Dummy() {}
+    __device__ ~Dummy() {}
+  };
+  __global__ void log10_kernel_v1(double* x) { double result = log10(x); }
+  __global__ void log10_kernel_v2(Dummy x) { double result = log10(x); }
+  __global__ void log10f_kernel_v1(float* x) { float result = log10f(x); }
+  __global__ void log10f_kernel_v2(Dummy x) { float result = log10f(x); }
+)"};
+
+static constexpr auto kLog1p{R"(
+  class Dummy {
+   public:
+    __device__ Dummy() {}
+    __device__ ~Dummy() {}
+  };
+  __global__ void log1p_kernel_v1(double* x) { double result = log1p(x); }
+  __global__ void log1p_kernel_v2(Dummy x) { double result = log1p(x); }
+  __global__ void log1pf_kernel_v1(float* x) { float result = log1pf(x); }
+  __global__ void log1pf_kernel_v2(Dummy x) { float result = log1pf(x); }
+)"};
+
+static constexpr auto kLogb{R"(
+  class Dummy {
+   public:
+    __device__ Dummy() {}
+    __device__ ~Dummy() {}
+  };
+  __global__ void logb_kernel_v1(double* x) { double result = logb(x); }
+  __global__ void logb_kernel_v2(Dummy x) { double result = logb(x); }
+  __global__ void logbf_kernel_v1(float* x) { float result = logbf(x); }
+  __global__ void logbf_kernel_v2(Dummy x) { float result = logbf(x); }
+)"};
+
+static constexpr auto kIlogb{R"(
+  class Dummy {
+   public:
+    __device__ Dummy() {}
+    __device__ ~Dummy() {}
+  };
+  __global__ void ilogb_kernel_v1(double* x) { double result = ilogb(x); }
+  __global__ void ilogb_kernel_v2(Dummy x) { double result = ilogb(x); }
+  __global__ void ilogbf_kernel_v1(float* x) { float result = ilogbf(x); }
+  __global__ void ilogbf_kernel_v2(Dummy x) { float result = ilogbf(x); }
+)"};
@@ -139,4 +139,6 @@ void TernaryFloatingPointTest(kernel_sig<T, TArg, TArg, TArg> kernel, ref_sig<RT
                                                                                                   \
    TernaryFloatingPointTest(kern_name##_kernel<TestType>, ref,                                    \
                            ULPValidatorBuilderFactory<TestType>(ulp));                            \
+                                                                                                   \
  }
+