diff --git a/projects/hip-tests/catch/unit/math/CMakeLists.txt b/projects/hip-tests/catch/unit/math/CMakeLists.txt index 2ebf11063d..380e1c0d4a 100644 --- a/projects/hip-tests/catch/unit/math/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/math/CMakeLists.txt @@ -20,6 +20,7 @@ set(TEST_SRC trig_funcs.cc + misc_funcs.cc ) hip_add_exe_to_target(NAME MathsTest @@ -34,4 +35,8 @@ add_test(NAME Unit_Device_Single_Precision_Trig_Functions_Negative add_test(NAME Unit_Device_Double_Precision_Trig_Functions_Negative COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} - trig_double_precision_negative_kernels.cc 66) \ No newline at end of file + trig_double_precision_negative_kernels.cc 66) +add_test(NAME Unit_Device_Misc_Functions_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + misc_negative_kernels.cc 76) diff --git a/projects/hip-tests/catch/unit/math/misc_funcs.cc b/projects/hip-tests/catch/unit/math/misc_funcs.cc new file mode 100644 index 0000000000..35e21fb26e --- /dev/null +++ b/projects/hip-tests/catch/unit/math/misc_funcs.cc @@ -0,0 +1,96 @@ +/* +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 "misc_negative_kernels_rtc.hh" + +#include "unary_common.hh" +#include "binary_common.hh" +#include "ternary_common.hh" + +MATH_UNARY_WITHIN_ULP_TEST_DEF(fabs, std::fabs, 0, 0) +TEST_CASE("Unit_Device_fabs_fabsf_Negative_RTC") { NegativeTestRTCWrapper<4>(kFabs); } + +MATH_BINARY_WITHIN_ULP_TEST_DEF(copysign, std::copysign, 0, 0) +TEST_CASE("Unit_Device_copysign_copysignf_Negative_RTC") { NegativeTestRTCWrapper<8>(kCopySign); } + +MATH_BINARY_WITHIN_ULP_TEST_DEF(fmax, std::fmax, 0, 0) +TEST_CASE("Unit_Device_fmax_fmaxf_Negative_RTC") { NegativeTestRTCWrapper<8>(kFmax); } + +MATH_BINARY_WITHIN_ULP_TEST_DEF(fmin, std::fmin, 0, 0) +TEST_CASE("Unit_Device_fmin_fminf_Negative_RTC") { NegativeTestRTCWrapper<8>(kFmin); } + +MATH_BINARY_WITHIN_ULP_TEST_DEF(nextafter, std::nextafter, 0, 0) +TEST_CASE("Unit_Device_nextafter_nextafterf_Negative_RTC") { + NegativeTestRTCWrapper<8>(kNextAfter); +} + +MATH_TERNARY_WITHIN_ULP_TEST_DEF(fma, std::fma, 0, 0) +TEST_CASE("Unit_Device_fma_fmaf_Negative_RTC") { NegativeTestRTCWrapper<12>(kFma); } + +__global__ void fdividef_kernel(float* const ys, const size_t num_xs, float* const x1s, + float* const x2s) { + const auto tid = cg::this_grid().thread_rank(); + const auto stride = cg::this_grid().size(); + + for (auto i = tid; i < num_xs; i += stride) { + ys[i] = fdividef(x1s[i], x2s[i]); + } +} + +TEST_CASE("Unit_Device_fdividef_Accuracy_Positive") { + double (*ref)(double, double) = [](double x1, double x2) { return x1 / x2; }; + BinaryFloatingPointTest(fdividef_kernel, ref, ULPValidatorBuilderFactory(0)); +} + +TEST_CASE("Unit_Device_fdividef_Negative_RTC") { NegativeTestRTCWrapper<4>(kFdividef); } + +#define MATH_BOOL_RETURNING_FUNCTION_TEST_DEF(kern_name, ref_func) \ + template \ + __global__ void kern_name##_kernel(bool* 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) { \ + ys[i] = kern_name(xs[i]); \ + } \ + } \ + \ + TEST_CASE("Unit_Device_" #kern_name "_Accuracy_Positive - float") { \ + bool (*ref)(double) = ref_func; \ + UnarySinglePrecisionTest(kern_name##_kernel, ref, EqValidatorBuilderFactory()); \ + } \ + \ + TEST_CASE("Unit_Device_" #kern_name "_Accuracy_Positive - double") { \ + bool (*ref)(long double) = ref_func; \ + UnaryDoublePrecisionTest(kern_name##_kernel, ref, EqValidatorBuilderFactory()); \ + } + +MATH_BOOL_RETURNING_FUNCTION_TEST_DEF(isfinite, std::isfinite) +TEST_CASE("Unit_Device_isfinite_Negative_RTC") { NegativeTestRTCWrapper<4>(kIsFinite); } + +MATH_BOOL_RETURNING_FUNCTION_TEST_DEF(isinf, std::isinf) +TEST_CASE("Unit_Device_isinf_Negative_RTC") { NegativeTestRTCWrapper<4>(kIsInf); } + +MATH_BOOL_RETURNING_FUNCTION_TEST_DEF(isnan, std::isnan) +TEST_CASE("Unit_Device_isnan_Negative_RTC") { NegativeTestRTCWrapper<4>(kIsNan); } + +MATH_BOOL_RETURNING_FUNCTION_TEST_DEF(signbit, std::signbit) +TEST_CASE("Unit_Device_signbit_Negative_RTC") { NegativeTestRTCWrapper<4>(kSignBit); } \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/math/misc_negative_kernels.cc b/projects/hip-tests/catch/unit/math/misc_negative_kernels.cc new file mode 100644 index 0000000000..761bd9cebf --- /dev/null +++ b/projects/hip-tests/catch/unit/math/misc_negative_kernels.cc @@ -0,0 +1,87 @@ +/* +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 + +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +#define MISC_UNARY_NEGATIVE_KERNELS(func_name) \ + __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); } \ + __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); } + +#define MISC_UNARY_BOOL_RET_NEGATIVE_KERNELS(func_name) \ + __global__ void func_name##_kernel_v1(float* x) { bool result = func_name(x); } \ + __global__ void func_name##_kernel_v2(Dummy x) { bool result = func_name(x); } \ + __global__ void func_name##_kernel_v3(double* x) { bool result = func_name(x); } \ + __global__ void func_name##_kernel_v4(Dummy x) { bool result = func_name(x); } + +#define MISC_BINARY_NEGATIVE_KERNELS(func_name) \ + __global__ void func_name##f_kernel_v1(float* x, float y) { float result = func_name##f(x, y); } \ + __global__ void func_name##f_kernel_v2(Dummy x, float y) { float result = func_name##f(x, y); } \ + __global__ void func_name##f_kernel_v3(float x, float* y) { float result = func_name##f(x, y); } \ + __global__ void func_name##f_kernel_v4(float x, Dummy y) { float result = func_name##f(x, y); } \ + __global__ void func_name##_kernel_v1(double* x, double y) { double result = func_name(x, y); } \ + __global__ void func_name##_kernel_v2(Dummy x, double y) { double result = func_name(x, y); } \ + __global__ void func_name##_kernel_v3(double x, double* y) { double result = func_name(x, y); } \ + __global__ void func_name##_kernel_v4(double x, Dummy y) { double result = func_name(x, y); } + +/*Expecting 4 errors*/ +MISC_UNARY_NEGATIVE_KERNELS(fabs) + +/*Expecting 8 errors per macro invocation - 40 total*/ +MISC_BINARY_NEGATIVE_KERNELS(copysign) +MISC_BINARY_NEGATIVE_KERNELS(fmax) +MISC_BINARY_NEGATIVE_KERNELS(fmin) +MISC_BINARY_NEGATIVE_KERNELS(nextafter) +MISC_BINARY_NEGATIVE_KERNELS(fma) + +/*Expecting 4 errors*/ +__global__ void fdividef_kernel_v1(float* x, float y) { float result = fdividef(x, y); } +__global__ void fdividef_kernel_v2(Dummy x, float y) { float result = fdivide(x); } +__global__ void fdividef_kernel_v3(float x, float* y) { float result = fdivide(x); } +__global__ void fdividef_kernel_v4(float x, Dummy y) { float result = fdivide(x); } + +/*Expecting 4 errors per macro invocation - 16 total*/ +MISC_UNARY_BOOL_RET_NEGATIVE_KERNELS(isfinite) +MISC_UNARY_BOOL_RET_NEGATIVE_KERNELS(isinf) +MISC_UNARY_BOOL_RET_NEGATIVE_KERNELS(isnan) +MISC_UNARY_BOOL_RET_NEGATIVE_KERNELS(signbit) + +/*Expecting 12 errors*/ +__global__ void fmaf_kernel_v1(float* x, float y, float z) { float result = fmaf(x, y, z); } +__global__ void fmaf_kernel_v2(Dummy x, float y, float z) { float result = fmaf(x, y, z); } +__global__ void fmaf_kernel_v3(float x, float* y, float z) { float result = fmaf(x, y, z); } +__global__ void fmaf_kernel_v4(float x, Dummy y, float z) { float result = fmaf(x, y, z); } +__global__ void fmaf_kernel_v5(float x, float y, float* z) { float result = fmaf(x, y, z); } +__global__ void fmaf_kernel_v6(float x, float y, Dummy z) { float result = fmaf(x, y, z); } +__global__ void fma_kernel_v1(double* x, double y, double z) { double result = fmaf(x, y, z); } +__global__ void fma_kernel_v2(Dummy x, double y, double z) { double result = fmaf(x, y, z); } +__global__ void fma_kernel_v3(double x, double* y, double z) { double result = fmaf(x, y, z); } +__global__ void fma_kernel_v4(double x, Dummy y, double z) { double result = fmaf(x, y, z); } +__global__ void fma_kernel_v5(double x, double y, double* z) { double result = fmaf(x, y, z); } +__global__ void fma_kernel_v6(double x, double y, Dummy z) { double result = fmaf(x, y, z); } \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/math/misc_negative_kernels_rtc.hh b/projects/hip-tests/catch/unit/math/misc_negative_kernels_rtc.hh new file mode 100644 index 0000000000..66521da090 --- /dev/null +++ b/projects/hip-tests/catch/unit/math/misc_negative_kernels_rtc.hh @@ -0,0 +1,177 @@ +/* +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. +*/ + + +static constexpr auto kFabs{R"( +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; +__global__ void fabsf_kernel_v1(float* x) { float result = fabsf(x); } +__global__ void fabsf_kernel_v2(Dummy x) { float result = fabsf(x); } +__global__ void fabs_kernel_v1(double* x) { double result = fabs(x); } +__global__ void fabs_kernel_v2(Dummy x) { double result = fabs(x); } +)"}; + +static constexpr auto kCopySign{R"( +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; +__global__ void copysignf_kernel_v1(float* x, float y) { float result = copysignf(x, y); } +__global__ void copysignf_kernel_v2(Dummy x, float y) { float result = copysignf(x, y); } +__global__ void copysignf_kernel_v3(float x, float* y) { float result = copysignf(x, y); } +__global__ void copysignf_kernel_v4(float x, Dummy y) { float result = copysignf(x, y); } +__global__ void copysign_kernel_v1(double* x, double y) { double result = copysign(x, y); } +__global__ void copysign_kernel_v2(Dummy x, double y) { double result = copysign(x, y); } +__global__ void copysign_kernel_v3(double x, double* y) { double result = copysign(x, y); } +__global__ void copysign_kernel_v4(double x, Dummy y) { double result = copysign(x, y); } +)"}; + +static constexpr auto kFmax{R"( +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; +__global__ void fmaxf_kernel_v1(float* x, float y) { float result = fmaxf(x, y); } +__global__ void fmaxf_kernel_v2(Dummy x, float y) { float result = fmaxf(x, y); } +__global__ void fmaxf_kernel_v3(float x, float* y) { float result = fmaxf(x, y); } +__global__ void fmaxf_kernel_v4(float x, Dummy y) { float result = fmaxf(x, y); } +__global__ void fmax_kernel_v1(double* x, double y) { double result = fmax(x, y); } +__global__ void fmax_kernel_v2(Dummy x, double y) { double result = fmax(x, y); } +__global__ void fmax_kernel_v3(double x, double* y) { double result = fmax(x, y); } +__global__ void fmax_kernel_v4(double x, Dummy y) { double result = fmax(x, y); } +)"}; + +static constexpr auto kFmin{R"( +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; +__global__ void fminf_kernel_v1(float* x, float y) { float result = fminf(x, y); } +__global__ void fminf_kernel_v2(Dummy x, float y) { float result = fminf(x, y); } +__global__ void fminf_kernel_v3(float x, float* y) { float result = fminf(x, y); } +__global__ void fminf_kernel_v4(float x, Dummy y) { float result = fminf(x, y); } +__global__ void fmin_kernel_v1(double* x, double y) { double result = fmin(x, y); } +__global__ void fmin_kernel_v2(Dummy x, double y) { double result = fmin(x, y); } +__global__ void fmin_kernel_v3(double x, double* y) { double result = fmin(x, y); } +__global__ void fmin_kernel_v4(double x, Dummy y) { double result = fmin(x, y); } +)"}; + +static constexpr auto kNextAfter{R"( +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; +__global__ void nextafterf_kernel_v1(float* x, float y) { float result = nextafterf(x, y); } +__global__ void nextafterf_kernel_v2(Dummy x, float y) { float result = nextafterf(x, y); } +__global__ void nextafterf_kernel_v3(float x, float* y) { float result = nextafterf(x, y); } +__global__ void nextafterf_kernel_v4(float x, Dummy y) { float result = nextafterf(x, y); } +__global__ void nextafter_kernel_v1(double* x, double y) { double result = nextafter(x, y); } +__global__ void nextafter_kernel_v2(Dummy x, double y) { double result = nextafter(x, y); } +__global__ void nextafter_kernel_v3(double x, double* y) { double result = nextafter(x, y); } +__global__ void nextafter_kernel_v4(double x, Dummy y) { double result = nextafter(x, y); } +)"}; + +static constexpr auto kFma{R"( +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; +__global__ void fmaf_kernel_v1(float* x, float y, float z) { float result = fmaf(x, y, z); } +__global__ void fmaf_kernel_v2(Dummy x, float y, float z) { float result = fmaf(x, y, z); } +__global__ void fmaf_kernel_v3(float x, float* y, float z) { float result = fmaf(x, y, z); } +__global__ void fmaf_kernel_v4(float x, Dummy y, float z) { float result = fmaf(x, y, z); } +__global__ void fmaf_kernel_v5(float x, float y, float* z) { float result = fmaf(x, y, z); } +__global__ void fmaf_kernel_v6(float x, float y, Dummy z) { float result = fmaf(x, y, z); } +__global__ void fma_kernel_v1(double* x, double y, double z) { double result = fmaf(x, y, z); } +__global__ void fma_kernel_v2(Dummy x, double y, double z) { double result = fmaf(x, y, z); } +__global__ void fma_kernel_v3(double x, double* y, double z) { double result = fmaf(x, y, z); } +__global__ void fma_kernel_v4(double x, Dummy y, double z) { double result = fmaf(x, y, z); } +__global__ void fma_kernel_v5(double x, double y, double* z) { double result = fmaf(x, y, z); } +__global__ void fma_kernel_v6(double x, double y, Dummy z) { double result = fmaf(x, y, z); } +)"}; + +static constexpr auto kFdividef{R"( +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; +__global__ void fdividef_kernel_v1(float* x, float y) { float result = fdividef(x, y); } +__global__ void fdividef_kernel_v2(Dummy x, float y) { float result = fdivide(x); } +__global__ void fdividef_kernel_v3(float x, float* y) { float result = fdivide(x); } +__global__ void fdividef_kernel_v4(float x, Dummy y) { float result = fdivide(x); } +)"}; + +static constexpr auto kIsFinite{R"( +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; +__global__ void isfinite_kernel_v1(float* x) { bool result = isfinite(x); } +__global__ void isfinite_kernel_v2(Dummy x) { bool result = isfinite(x); } +__global__ void isfinite_kernel_v3(double* x) { bool result = isfinite(x); } +__global__ void isfinite_kernel_v4(Dummy x) { bool result = isfinite(x); } +)"}; + +static constexpr auto kIsInf{R"( +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; +__global__ void isinf_kernel_v1(float* x) { bool result = isinf(x); } +__global__ void isinf_kernel_v2(Dummy x) { bool result = isinf(x); } +__global__ void isinf_kernel_v3(double* x) { bool result = isinf(x); } +__global__ void isinf_kernel_v4(Dummy x) { bool result = isinf(x); } +)"}; + +static constexpr auto kIsNan{R"( +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; +__global__ void isnan_kernel_v1(float* x) { bool result = isnan(x); } +__global__ void isnan_kernel_v2(Dummy x) { bool result = isnan(x); } +__global__ void isnan_kernel_v3(double* x) { bool result = isnan(x); } +__global__ void isnan_kernel_v4(Dummy x) { bool result = isnan(x); } +)"}; + +static constexpr auto kSignBit{R"( +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; +__global__ void signbit_kernel_v1(float* x) { bool result = signbit(x); } +__global__ void signbit_kernel_v2(Dummy x) { bool result = signbit(x); } +__global__ void signbit_kernel_v3(double* x) { bool result = signbit(x); } +__global__ void signbit_kernel_v4(Dummy x) { bool result = signbit(x); } +)"};