diff --git a/catch/include/hip_test_common.hh b/catch/include/hip_test_common.hh index c2d19650bf..21707f7615 100644 --- a/catch/include/hip_test_common.hh +++ b/catch/include/hip_test_common.hh @@ -178,7 +178,7 @@ static inline bool IsGfx11() { hipDeviceProp_t props{}; HIP_CHECK(hipGetDevice(&device)); HIP_CHECK(hipGetDeviceProperties(&props, device)); - // Get GCN Arch Name and compare to check if it is gfx11 + // Get GCN Arch Name and compare to check if it is gfx11 std::string arch = std::string(props.gcnArchName); auto pos = arch.find("gfx11"); if (pos != std::string::npos) @@ -186,7 +186,7 @@ static inline bool IsGfx11() { else return false; #else - std::cout<<"Have to be either Nvidia or AMD platform, asserting"<(kernel, numBlocks, numThreads, memPerBlock, stream, std::forward(packedArgs)...); #endif -HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipGetLastError()); } //--- diff --git a/catch/unit/math/CMakeLists.txt b/catch/unit/math/CMakeLists.txt index 380e1c0d4a..b0a2b5d00c 100644 --- a/catch/unit/math/CMakeLists.txt +++ b/catch/unit/math/CMakeLists.txt @@ -21,11 +21,19 @@ set(TEST_SRC trig_funcs.cc misc_funcs.cc + remainder_and_rounding_funcs.cc ) +if(HIP_PLATFORM MATCHES "nvidia") + set(LINKER_LIBS nvrtc) +elseif(HIP_PLATFORM MATCHES "amd") + set(LINKER_LIBS hiprtc) +endif() + hip_add_exe_to_target(NAME MathsTest TEST_SRC ${TEST_SRC} - TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC}) + TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC} + LINKER_LIBS ${LINKER_LIBS}) add_test(NAME Unit_Device_Single_Precision_Trig_Functions_Negative COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py @@ -40,3 +48,13 @@ 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) + +add_test(NAME Unit_Device_remainder_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + math_remainder_negative_kernels.cc 68) + +add_test(NAME Unit_Device_rounding_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + math_rounding_negative_kernels.cc 40) diff --git a/catch/unit/math/math_common.hh b/catch/unit/math/math_common.hh index 8b59558389..7ebc9b8f5d 100644 --- a/catch/unit/math/math_common.hh +++ b/catch/unit/math/math_common.hh @@ -216,7 +216,13 @@ template void NegativeTestRTCWrapper(const char* program_source) HIPRTC_CHECK( hiprtcCreateProgram(&program, program_source, "math_test_rtc.cc", 0, nullptr, nullptr)); +#if HT_AMD + std::string args = std::string("-ferror-limit=100"); + const char* options[] = {args.c_str()}; + hiprtcResult result{hiprtcCompileProgram(program, 1, options)}; +#else hiprtcResult result{hiprtcCompileProgram(program, 0, nullptr)}; +#endif // Get the compile log and count compiler error messages size_t log_size{}; diff --git a/catch/unit/math/math_remainder_negative_kernels.cc b/catch/unit/math/math_remainder_negative_kernels.cc new file mode 100644 index 0000000000..2ebd26516e --- /dev/null +++ b/catch/unit/math/math_remainder_negative_kernels.cc @@ -0,0 +1,113 @@ +/* +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 + +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +#define NEGATIVE_KERNELS_SHELL(func_name) \ + __global__ void func_name##_kernel_v1(double* x, double y) { auto result = func_name(x, y); } \ + __global__ void func_name##_kernel_v2(double x, double* y) { auto result = func_name(x, y); } \ + __global__ void func_name##_kernel_v3(Dummy x, double y) { auto result = func_name(x, y); } \ + __global__ void func_name##_kernel_v4(double x, Dummy y) { auto result = func_name(x, y); } \ + __global__ void func_name##f_kernel_v1(float* x, float y) { auto result = func_name##f(x, y); } \ + __global__ void func_name##f_kernel_v2(float x, float* y) { auto result = func_name##f(x, y); } \ + __global__ void func_name##f_kernel_v3(Dummy x, float y) { auto result = func_name##f(x, y); } \ + __global__ void func_name##f_kernel_v4(float x, Dummy y) { auto result = func_name##f(x, y); } + +NEGATIVE_KERNELS_SHELL(fmod) +NEGATIVE_KERNELS_SHELL(remainder) + +__global__ void remquo_kernel_v1(double* x, double y, int* quo) { auto result = remquo(x, y, quo); } +__global__ void remquo_kernel_v2(Dummy x, double y, int* quo) { auto result = remquo(x, y, quo); } +__global__ void remquo_kernel_v3(double x, double* y, int* quo) { auto result = remquo(x, y, quo); } +__global__ void remquo_kernel_v4(double x, Dummy y, int* quo) { auto result = remquo(x, y, quo); } +__global__ void remquo_kernel_v5(double x, double y, char* quo) { auto result = remquo(x, y, quo); } +__global__ void remquo_kernel_v6(double x, double y, short* quo) { + auto result = remquo(x, y, quo); +} +__global__ void remquo_kernel_v7(double x, double y, long* quo) { auto result = remquo(x, y, quo); } +__global__ void remquo_kernel_v8(double x, double y, long long* quo) { + auto result = remquo(x, y, quo); +} +__global__ void remquo_kernel_v9(double x, double y, float* quo) { + auto result = remquo(x, y, quo); +} +__global__ void remquo_kernel_v10(double x, double y, double* quo) { + auto result = remquo(x, y, quo); +} +__global__ void remquo_kernel_v11(double x, double y, Dummy* quo) { + auto result = remquo(x, y, quo); +} +__global__ void remquo_kernel_v12(double x, double y, const int* quo) { + auto result = remquo(x, y, quo); +} + +__global__ void remquof_kernel_v1(float* x, float y, int* quo) { auto result = remquof(x, y, quo); } +__global__ void remquof_kernel_v2(Dummy x, float y, int* quo) { auto result = remquof(x, y, quo); } +__global__ void remquof_kernel_v3(float x, float* y, int* quo) { auto result = remquof(x, y, quo); } +__global__ void remquof_kernel_v4(float x, Dummy y, int* quo) { auto result = remquof(x, y, quo); } +__global__ void remquof_kernel_v5(float x, float y, char* quo) { auto result = remquof(x, y, quo); } +__global__ void remquof_kernel_v6(float x, float y, short* quo) { + auto result = remquof(x, y, quo); +} +__global__ void remquof_kernel_v7(float x, float y, long* quo) { auto result = remquof(x, y, quo); } +__global__ void remquof_kernel_v8(float x, float y, long long* quo) { + auto result = remquof(x, y, quo); +} +__global__ void remquof_kernel_v9(float x, float y, float* quo) { + auto result = remquof(x, y, quo); +} +__global__ void remquof_kernel_v10(float x, float y, double* quo) { + auto result = remquof(x, y, quo); +} +__global__ void remquof_kernel_v11(float x, float y, Dummy* quo) { + auto result = remquof(x, y, quo); +} +__global__ void remquof_kernel_v12(float x, float y, const int* quo) { + auto result = remquof(x, y, quo); +} + +__global__ void modf_kernel_v1(double* x, double* iptr) { auto result = modf(x, iptr); } +__global__ void modf_kernel_v2(Dummy x, double* iptr) { auto result = modf(x, iptr); } +__global__ void modf_kernel_v3(double x, int* iptr) { auto result = modf(x, iptr); } +__global__ void modf_kernel_v4(double x, char* iptr) { auto result = modf(x, iptr); } +__global__ void modf_kernel_v5(double x, short* iptr) { auto result = modf(x, iptr); } +__global__ void modf_kernel_v6(double x, long* iptr) { auto result = modf(x, iptr); } +__global__ void modf_kernel_v7(double x, long long* iptr) { auto result = modf(x, iptr); } +__global__ void modf_kernel_v8(double x, float* iptr) { auto result = modf(x, iptr); } +__global__ void modf_kernel_v9(double x, Dummy* iptr) { auto result = modf(x, iptr); } +__global__ void modf_kernel_v10(double x, const double* iptr) { auto result = modf(x, iptr); } + +__global__ void modff_kernel_v1(float* x, float* iptr) { auto result = modff(x, iptr); } +__global__ void modff_kernel_v2(Dummy x, float* iptr) { auto result = modff(x, iptr); } +__global__ void modff_kernel_v3(float x, int* iptr) { auto result = modff(x, iptr); } +__global__ void modff_kernel_v4(float x, char* iptr) { auto result = modff(x, iptr); } +__global__ void modff_kernel_v5(float x, short* iptr) { auto result = modff(x, iptr); } +__global__ void modff_kernel_v6(float x, long* iptr) { auto result = modff(x, iptr); } +__global__ void modff_kernel_v7(float x, long long* iptr) { auto result = modff(x, iptr); } +__global__ void modff_kernel_v8(float x, double* iptr) { auto result = modff(x, iptr); } +__global__ void modff_kernel_v9(float x, Dummy* iptr) { auto result = modff(x, iptr); } +__global__ void modff_kernel_v10(float x, const float* iptr) { auto result = modff(x, iptr); } + +NEGATIVE_KERNELS_SHELL(fdim) diff --git a/catch/unit/math/math_remainder_rounding_negative_kernels_rtc.hh b/catch/unit/math/math_remainder_rounding_negative_kernels_rtc.hh new file mode 100644 index 0000000000..e67a6d0092 --- /dev/null +++ b/catch/unit/math/math_remainder_rounding_negative_kernels_rtc.hh @@ -0,0 +1,276 @@ +/* +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 remainder and rounding negative Test Cases that are using RTC. +*/ + +static constexpr auto kTrunc{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void trunc_kernel_v1(double* x) { auto result = trunc(x); } + __global__ void trunc_kernel_v2(Dummy x) { auto result = trunc(x); } + __global__ void truncf_kernel_v1(float* x) { auto result = truncf(x); } + __global__ void truncf_kernel_v2(Dummy x) { auto result = truncf(x); } +)"}; + +static constexpr auto kRound{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void round_kernel_v1(double* x) { auto result = round(x); } + __global__ void round_kernel_v2(Dummy x) { auto result = round(x); } + __global__ void roundf_kernel_v1(float* x) { auto result = roundf(x); } + __global__ void roundf_kernel_v2(Dummy x) { auto result = roundf(x); } +)"}; + +static constexpr auto kRint{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void rint_kernel_v1(double* x) { auto result = rint(x); } + __global__ void rint_kernel_v2(Dummy x) { auto result = rint(x); } + __global__ void rintf_kernel_v1(float* x) { auto result = rintf(x); } + __global__ void rintf_kernel_v2(Dummy x) { auto result = rintf(x); } +)"}; + +static constexpr auto kNearbyint{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void nearbyint_kernel_v1(double* x) { auto result = nearbyint(x); } + __global__ void nearbyint_kernel_v2(Dummy x) { auto result = nearbyint(x); } + __global__ void nearbyintf_kernel_v1(float* x) { auto result = nearbyintf(x); } + __global__ void nearbyintf_kernel_v2(Dummy x) { auto result = nearbyintf(x); } +)"}; + +static constexpr auto kCeil{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void ceil_kernel_v1(double* x) { auto result = ceil(x); } + __global__ void ceil_kernel_v2(Dummy x) { auto result = ceil(x); } + __global__ void ceilf_kernel_v1(float* x) { auto result = ceilf(x); } + __global__ void ceilf_kernel_v2(Dummy x) { auto result = ceilf(x); } +)"}; + +static constexpr auto kFloor{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void floor_kernel_v1(double* x) { auto result = floor(x); } + __global__ void floor_kernel_v2(Dummy x) { auto result = floor(x); } + __global__ void floorf_kernel_v1(float* x) { auto result = floorf(x); } + __global__ void floorf_kernel_v2(Dummy x) { auto result = floorf(x); } +)"}; + +static constexpr auto kLrint{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void lrint_kernel_v1(double* x) { auto result = lrint(x); } + __global__ void lrint_kernel_v2(Dummy x) { auto result = lrint(x); } + __global__ void lrintf_kernel_v1(float* x) { auto result = lrintf(x); } + __global__ void lrintf_kernel_v2(Dummy x) { auto result = lrintf(x); } +)"}; + +static constexpr auto kLround{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void lround_kernel_v1(double* x) { auto result = lround(x); } + __global__ void lround_kernel_v2(Dummy x) { auto result = lround(x); } + __global__ void lroundf_kernel_v1(float* x) { auto result = lroundf(x); } + __global__ void lroundf_kernel_v2(Dummy x) { auto result = lroundf(x); } +)"}; + +static constexpr auto kLlrint{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void llrint_kernel_v1(double* x) { auto result = llrint(x); } + __global__ void llrint_kernel_v2(Dummy x) { auto result = llrint(x); } + __global__ void llrintf_kernel_v1(float* x) { auto result = llrintf(x); } + __global__ void llrintf_kernel_v2(Dummy x) { auto result = llrintf(x); } +)"}; + +static constexpr auto kLlround{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void llround_kernel_v1(double* x) { auto result = llround(x); } + __global__ void llround_kernel_v2(Dummy x) { auto result = llround(x); } + __global__ void llroundf_kernel_v1(float* x) { auto result = llroundf(x); } + __global__ void llroundf_kernel_v2(Dummy x) { auto result = llroundf(x); } +)"}; + +static constexpr auto kFmod{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void fmod_kernel_v1(double* x, double y) { auto result = fmod(x, y); } + __global__ void fmod_kernel_v2(double x, double* y) { auto result = fmod(x, y); } + __global__ void fmod_kernel_v3(Dummy x, double y) { auto result = fmod(x, y); } + __global__ void fmod_kernel_v4(double x, Dummy y) { auto result = fmod(x, y); } + __global__ void fmodf_kernel_v1(float* x, float y) { auto result = fmodf(x, y); } + __global__ void fmodf_kernel_v2(float x, float* y) { auto result = fmodf(x, y); } + __global__ void fmodf_kernel_v3(Dummy x, float y) { auto result = fmodf(x, y); } + __global__ void fmodf_kernel_v4(float x, Dummy y) { auto result = fmodf(x, y); } +)"}; + +static constexpr auto kRemainder{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void remainder_kernel_v1(double* x, double y) { auto result = remainder(x, y); } + __global__ void remainder_kernel_v2(double x, double* y) { auto result = remainder(x, y); } + __global__ void remainder_kernel_v3(Dummy x, double y) { auto result = remainder(x, y); } + __global__ void remainder_kernel_v4(double x, Dummy y) { auto result = remainder(x, y); } + __global__ void remainderf_kernel_v1(float* x, float y) { auto result = remainderf(x, y); } + __global__ void remainderf_kernel_v2(float x, float* y) { auto result = remainderf(x, y); } + __global__ void remainderf_kernel_v3(Dummy x, float y) { auto result = remainderf(x, y); } + __global__ void remainderf_kernel_v4(float x, Dummy y) { auto result = remainderf(x, y); } +)"}; + +static constexpr auto kRemquo{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void remquo_kernel_v1(double* x, double y, int* quo) { auto result = remquo(x, y, quo); } + __global__ void remquo_kernel_v2(Dummy x, double y, int* quo) { auto result = remquo(x, y, quo); } + __global__ void remquo_kernel_v3(double x, double* y, int* quo) { auto result = remquo(x, y, quo); } + __global__ void remquo_kernel_v4(double x, Dummy y, int* quo) { auto result = remquo(x, y, quo); } + __global__ void remquo_kernel_v5(double x, double y, char* quo) { auto result = remquo(x, y, quo); } + __global__ void remquo_kernel_v6(double x, double y, short* quo) { + auto result = remquo(x, y, quo); + } + __global__ void remquo_kernel_v7(double x, double y, long* quo) { auto result = remquo(x, y, quo); } + __global__ void remquo_kernel_v8(double x, double y, long long* quo) { + auto result = remquo(x, y, quo); + } + __global__ void remquo_kernel_v9(double x, double y, float* quo) { + auto result = remquo(x, y, quo); + } + __global__ void remquo_kernel_v10(double x, double y, double* quo) { + auto result = remquo(x, y, quo); + } + __global__ void remquo_kernel_v11(double x, double y, Dummy* quo) { + auto result = remquo(x, y, quo); + } + __global__ void remquo_kernel_v12(double x, double y, const int* quo) { + auto result = remquo(x, y, quo); + } + __global__ void remquof_kernel_v1(float* x, float y, int* quo) { auto result = remquof(x, y, quo); } + __global__ void remquof_kernel_v2(Dummy x, float y, int* quo) { auto result = remquof(x, y, quo); } + __global__ void remquof_kernel_v3(float x, float* y, int* quo) { auto result = remquof(x, y, quo); } + __global__ void remquof_kernel_v4(float x, Dummy y, int* quo) { auto result = remquof(x, y, quo); } + __global__ void remquof_kernel_v5(float x, float y, char* quo) { auto result = remquof(x, y, quo); } + __global__ void remquof_kernel_v6(float x, float y, short* quo) { + auto result = remquof(x, y, quo); + } + __global__ void remquof_kernel_v7(float x, float y, long* quo) { auto result = remquof(x, y, quo); } + __global__ void remquof_kernel_v8(float x, float y, long long* quo) { + auto result = remquof(x, y, quo); + } + __global__ void remquof_kernel_v9(float x, float y, float* quo) { + auto result = remquof(x, y, quo); + } + __global__ void remquof_kernel_v10(float x, float y, double* quo) { + auto result = remquof(x, y, quo); + } + __global__ void remquof_kernel_v11(float x, float y, Dummy* quo) { + auto result = remquof(x, y, quo); + } + __global__ void remquof_kernel_v12(float x, float y, const int* quo) { + auto result = remquof(x, y, quo); + } +)"}; + +static constexpr auto kModf{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void modf_kernel_v1(double* x, double* iptr) { auto result = modf(x, iptr); } + __global__ void modf_kernel_v2(Dummy x, double* iptr) { auto result = modf(x, iptr); } + __global__ void modf_kernel_v3(double x, int* iptr) { auto result = modf(x, iptr); } + __global__ void modf_kernel_v4(double x, char* iptr) { auto result = modf(x, iptr); } + __global__ void modf_kernel_v5(double x, short* iptr) { auto result = modf(x, iptr); } + __global__ void modf_kernel_v6(double x, long* iptr) { auto result = modf(x, iptr); } + __global__ void modf_kernel_v7(double x, long long* iptr) { auto result = modf(x, iptr); } + __global__ void modf_kernel_v8(double x, float* iptr) { auto result = modf(x, iptr); } + __global__ void modf_kernel_v9(double x, Dummy* iptr) { auto result = modf(x, iptr); } + __global__ void modf_kernel_v10(double x, const double* iptr) { auto result = modf(x, iptr); } + __global__ void modff_kernel_v1(float* x, float* iptr) { auto result = modff(x, iptr); } + __global__ void modff_kernel_v2(Dummy x, float* iptr) { auto result = modff(x, iptr); } + __global__ void modff_kernel_v3(float x, int* iptr) { auto result = modff(x, iptr); } + __global__ void modff_kernel_v4(float x, char* iptr) { auto result = modff(x, iptr); } + __global__ void modff_kernel_v5(float x, short* iptr) { auto result = modff(x, iptr); } + __global__ void modff_kernel_v6(float x, long* iptr) { auto result = modff(x, iptr); } + __global__ void modff_kernel_v7(float x, long long* iptr) { auto result = modff(x, iptr); } + __global__ void modff_kernel_v8(float x, double* iptr) { auto result = modff(x, iptr); } + __global__ void modff_kernel_v9(float x, Dummy* iptr) { auto result = modff(x, iptr); } + __global__ void modff_kernel_v10(float x, const float* iptr) { auto result = modff(x, iptr); } +)"}; + +static constexpr auto kFdim{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void fdim_kernel_v1(double* x, double y) { auto result = fdim(x, y); } + __global__ void fdim_kernel_v2(double x, double* y) { auto result = fdim(x, y); } + __global__ void fdim_kernel_v3(Dummy x, double y) { auto result = fdim(x, y); } + __global__ void fdim_kernel_v4(double x, Dummy y) { auto result = fdim(x, y); } + __global__ void fdimf_kernel_v1(float* x, float y) { auto result = fdimf(x, y); } + __global__ void fdimf_kernel_v2(float x, float* y) { auto result = fdimf(x, y); } + __global__ void fdimf_kernel_v3(Dummy x, float y) { auto result = fdimf(x, y); } + __global__ void fdimf_kernel_v4(float x, Dummy y) { auto result = fdimf(x, y); } +)"}; diff --git a/catch/unit/math/math_rounding_negative_kernels.cc b/catch/unit/math/math_rounding_negative_kernels.cc new file mode 100644 index 0000000000..857f50d5dd --- /dev/null +++ b/catch/unit/math/math_rounding_negative_kernels.cc @@ -0,0 +1,43 @@ +/* +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 + +class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} +}; + +#define NEGATIVE_KERNELS_SHELL(func_name) \ + __global__ void func_name##_kernel_v1(double* x) { auto result = func_name(x); } \ + __global__ void func_name##_kernel_v2(Dummy x) { auto result = func_name(x); } \ + __global__ void func_name##f_kernel_v1(float* x) { auto result = func_name##f(x); } \ + __global__ void func_name##f_kernel_v2(Dummy x) { auto result = func_name##f(x); } + +NEGATIVE_KERNELS_SHELL(trunc) +NEGATIVE_KERNELS_SHELL(round) +NEGATIVE_KERNELS_SHELL(rint) +NEGATIVE_KERNELS_SHELL(nearbyint) +NEGATIVE_KERNELS_SHELL(ceil) +NEGATIVE_KERNELS_SHELL(floor) +NEGATIVE_KERNELS_SHELL(lrint) +NEGATIVE_KERNELS_SHELL(lround) +NEGATIVE_KERNELS_SHELL(llrint) +NEGATIVE_KERNELS_SHELL(llround) diff --git a/catch/unit/math/remainder_and_rounding_funcs.cc b/catch/unit/math/remainder_and_rounding_funcs.cc new file mode 100644 index 0000000000..635a68a459 --- /dev/null +++ b/catch/unit/math/remainder_and_rounding_funcs.cc @@ -0,0 +1,153 @@ +/* +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 "binary_common.hh" +#include "math_remainder_rounding_negative_kernels_rtc.hh" + +MATH_BINARY_WITHIN_ULP_TEST_DEF(fmod, std::fmod, 0, 0) +TEST_CASE("Unit_Device_fmod_fmodf_Negative_RTC") { NegativeTestRTCWrapper<8>(kFmod); } + +MATH_BINARY_WITHIN_ULP_TEST_DEF(remainder, std::remainder, 0, 0) +TEST_CASE("Unit_Device_remainder_remainder_Negative_RTC") { NegativeTestRTCWrapper<8>(kRemainder); } + +MATH_BINARY_WITHIN_ULP_TEST_DEF(fdim, std::fdim, 0, 0) +TEST_CASE("Unit_Device_fdim_fdimf_Negative_RTC") { NegativeTestRTCWrapper<8>(kFdim); } + +MATH_UNARY_WITHIN_ULP_TEST_DEF(trunc, std::trunc, 0, 0) +TEST_CASE("Unit_Device_trunc_truncf_Negative_RTC") { NegativeTestRTCWrapper<4>(kTrunc); } + +MATH_UNARY_WITHIN_ULP_TEST_DEF(round, std::round, 0, 0) +TEST_CASE("Unit_Device_round_roundf_Negative_RTC") { NegativeTestRTCWrapper<4>(kRound); } + +MATH_UNARY_WITHIN_ULP_TEST_DEF(rint, std::rint, 0, 0) +TEST_CASE("Unit_Device_rint_rintf_Negative_RTC") { NegativeTestRTCWrapper<4>(kRint); } + +MATH_UNARY_WITHIN_ULP_TEST_DEF(nearbyint, std::nearbyint, 0, 0) +TEST_CASE("Unit_Device_nearbyint_nearbyintf_Negative_RTC") { + NegativeTestRTCWrapper<4>(kNearbyint); +} + +MATH_UNARY_WITHIN_ULP_TEST_DEF(ceil, std::ceil, 0, 0) +TEST_CASE("Unit_Device_ceil_ceilf_Negative_RTC") { NegativeTestRTCWrapper<4>(kCeil); } + +MATH_UNARY_WITHIN_ULP_TEST_DEF(floor, std::floor, 0, 0) +TEST_CASE("Unit_Device_floor_floorf_Negative_RTC") { NegativeTestRTCWrapper<4>(kFloor); } + + +#define LONG_CONVERSION_FUNCTION_TEST_DEF(kern_name, ref_func, lt) \ + MATH_UNARY_KERNEL_DEF(kern_name) \ + \ + TEST_CASE("Unit_Device_" #kern_name "_Accuracy_Positive - float") { \ + lt (*ref)(double) = ref_func; \ + UnarySinglePrecisionRangeTest(kern_name##_kernel, ref, \ + EqValidatorBuilderFactory(), \ + static_cast(std::numeric_limits::lowest()), \ + static_cast(std::numeric_limits::max())); \ + } \ + \ + TEST_CASE("Unit_Device_" #kern_name "_Accuracy_Positive - double") { \ + lt (*ref)(long double) = ref_func; \ + UnaryDoublePrecisionBruteForceTest(kern_name##_kernel, ref, \ + EqValidatorBuilderFactory(), \ + static_cast(std::numeric_limits::lowest()), \ + static_cast(std::numeric_limits::max())); \ + } + +LONG_CONVERSION_FUNCTION_TEST_DEF(lrint, std::lrint, long) +TEST_CASE("Unit_Device_lrint_lrintf_Negative_RTC") { NegativeTestRTCWrapper<4>(kLrint); } + +LONG_CONVERSION_FUNCTION_TEST_DEF(lround, std::lround, long) +TEST_CASE("Unit_Device_lround_lroundf_Negative_RTC") { NegativeTestRTCWrapper<4>(kLround); } + +LONG_CONVERSION_FUNCTION_TEST_DEF(llrint, std::llrint, long long) +TEST_CASE("Unit_Device_llrint_llrintf_Negative_RTC") { NegativeTestRTCWrapper<4>(kLlrint); } + +LONG_CONVERSION_FUNCTION_TEST_DEF(llround, std::llround, long long) +TEST_CASE("Unit_Device_llround_llroundf_Negative_RTC") { NegativeTestRTCWrapper<4>(kLlround); } + + +template +__global__ void remquo_kernel(std::pair* const ys, const size_t num_xs, T* const x1s, + T* 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) { + if constexpr (std::is_same_v) { + ys[i].first = remquof(x1s[i], x2s[i], &ys[i].second); + } else if constexpr (std::is_same_v) { + ys[i].first = remquo(x1s[i], x2s[i], &ys[i].second); + } + } +} + +template std::pair remquo_wrapper(T x1, T x2) { + std::pair ret; + ret.first = std::remquo(x1, x2, &ret.second); + return ret; +} + +TEMPLATE_TEST_CASE("Unit_Device_remquo_Accuracy_Positive", "", float, double) { + using RT = RefType_t; + std::pair (*ref)(RT, RT) = remquo_wrapper; + const auto ulp_builder = ULPValidatorBuilderFactory(0); + const auto eq_builder = EqValidatorBuilderFactory(); + + BinaryFloatingPointTest(remquo_kernel, ref, + PairValidatorBuilderFactory(ulp_builder, eq_builder)); +} + +TEST_CASE("Unit_Device_remquo_remquof_Negative_RTC") { NegativeTestRTCWrapper<24>(kRemquo); } + +template +__global__ void modf_kernel(std::pair* 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) { + ys[i].first = modff(xs[i], &ys[i].second); + } else if constexpr (std::is_same_v) { + ys[i].first = modf(xs[i], &ys[i].second); + } + } +} + +template std::pair modf_wrapper(T x) { + std::pair ret; + ret.first = std::modf(x, &ret.second); + return ret; +} + +TEST_CASE("Unit_Device_modf_Accuracy_Positive - float") { + UnarySinglePrecisionTest( + modf_kernel, modf_wrapper, + PairValidatorBuilderFactory(ULPValidatorBuilderFactory(0))); +} + +TEST_CASE("Unit_Device_modf_Accuracy_Positive - double") { + UnaryDoublePrecisionTest( + modf_kernel, modf_wrapper, + PairValidatorBuilderFactory(ULPValidatorBuilderFactory(0))); +} + +TEST_CASE("Unit_Device_modf_modff_Negative_RTC") { NegativeTestRTCWrapper<20>(kModf); } diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index fda74f5b2e..d99cf33b32 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -83,17 +83,7 @@ if(HIP_PLATFORM MATCHES "amd") hipExtMallocWithFlags.cc hipMallocMngdMultiThread.cc hipArray.cc - hipMemVmm.cc - hipMemCreate.cc - hipMemMap.cc - hipMemGetAllocationGranularity.cc - hipMemSetGetAccess.cc - hipMemRetainAllocationHandle.cc - hipMemUnmap.cc - hipMemAddressFree.cc - hipMemAddressReserve.cc - hipMemRelease.cc - hipMemGetAllocationPropertiesFromHandle.cc) + hipMemVmm.cc) else() set(TEST_SRC ${TEST_SRC} hipGetSymbolSizeAddress.cc) endif() @@ -138,8 +128,10 @@ set(TEST_SRC hipMemsetFunctional.cc hipMalloc.cc hipMallocPitch.cc + hipMallocArray.cc hipMalloc3D.cc hipMalloc3DArray.cc + hipArrayCreate.cc hipArray3DCreate.cc hipDrvMemcpy3D.cc hipDrvMemcpy3D_old.cc @@ -165,52 +157,6 @@ set(TEST_SRC hipMallocMipmappedArray.cc hipFreeMipmappedArray.cc) -set(NOT_FOR_MI200_AND_ABOVE_TEST hipMallocArray.cc hipArrayCreate.cc) # tests not for MI200+ -set(MI200_AND_ABOVE_TARGETS gfx90a gfx940 gfx941 gfx942) -function(CheckRejectedArchs OFFLOAD_ARCH_STR_LOCAL) - set(ARCH_CHECK -1 PARENT_SCOPE) - string(REGEX MATCHALL "--offload-arch=gfx[0-9a-z]+" OFFLOAD_ARCH_LIST ${OFFLOAD_ARCH_STR_LOCAL}) - foreach(OFFLOAD_ARCH IN LISTS OFFLOAD_ARCH_LIST) - string(REGEX MATCHALL "--offload-arch=(gfx[0-9a-z]+)" matches ${OFFLOAD_ARCH}) - if (CMAKE_MATCH_COUNT EQUAL 1) - if (CMAKE_MATCH_1 IN_LIST MI200_AND_ABOVE_TARGETS) - set(ARCH_CHECK 1 PARENT_SCOPE) - endif() # CMAKE_MATCH_1 - endif() # CMAKE_MATCH_COUNT - endforeach() # OFFLOAD_ARCH_LIST -endfunction() # CheckAcceptedArchs - -if(HIP_PLATFORM MATCHES "amd") - if (DEFINED OFFLOAD_ARCH_STR) - CheckRejectedArchs(${OFFLOAD_ARCH_STR}) - elseif(DEFINED $ENV{HCC_AMDGPU_TARGET}) - CheckRejectedArchs($ENV{HCC_AMDGPU_TARGET}) - else() - set(ARCH_CHECK -1) - endif() - if(${ARCH_CHECK} EQUAL -1) - message(STATUS "Adding test: ${NOT_FOR_MI200_AND_ABOVE_TEST}") - set(TEST_SRC ${TEST_SRC} ${NOT_FOR_MI200_AND_ABOVE_TEST}) - endif() -else() - set(TEST_SRC ${TEST_SRC} ${NOT_FOR_MI200_AND_ABOVE_TEST}) -endif() - hip_add_exe_to_target(NAME MemoryTest2 TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC}) - -if(HIP_PLATFORM MATCHES "amd") - set(TEST_SRC - hipSVMTestByteGranularity.cpp - hipSVMTestFineGrainMemoryConsistency.cpp - hipSVMTestFineGrainSyncBuffers.cpp - hipSVMTestSharedAddressSpaceFineGrain.cpp - ) - - hip_add_exe_to_target(NAME SVMAtomicTest - TEST_SRC ${TEST_SRC} - TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC}) - - add_dependencies(build_tests hipHostRegisterPerf) -endif()