EXSWHTEC-289 - Implement tests for remainder and rounding device math functions #234
Change-Id: I8413cfeb0cbf32e2e8369c5b1527c9794a595688
This commit is contained in:
Dino Music
Rakesh Roy
@@ -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"<<std::endl;
|
||||
std::cout << "Have to be either Nvidia or AMD platform, asserting" << std::endl;
|
||||
assert(false);
|
||||
#endif
|
||||
}
|
||||
@@ -321,7 +321,7 @@ void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerB
|
||||
launchRTCKernel<Typenames...>(kernel, numBlocks, numThreads, memPerBlock, stream,
|
||||
std::forward<Args>(packedArgs)...);
|
||||
#endif
|
||||
HIP_CHECK(hipGetLastError());
|
||||
HIP_CHECK(hipGetLastError());
|
||||
}
|
||||
|
||||
//---
|
||||
|
||||
@@ -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)
|
||||
|
||||
@@ -216,7 +216,13 @@ template <int error_num> 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{};
|
||||
|
||||
@@ -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 <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 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)
|
||||
@@ -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); }
|
||||
)"};
|
||||
@@ -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 <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) { 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)
|
||||
@@ -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<float, lt>, ref, \
|
||||
EqValidatorBuilderFactory<lt>(), \
|
||||
static_cast<float>(std::numeric_limits<lt>::lowest()), \
|
||||
static_cast<float>(std::numeric_limits<lt>::max())); \
|
||||
} \
|
||||
\
|
||||
TEST_CASE("Unit_Device_" #kern_name "_Accuracy_Positive - double") { \
|
||||
lt (*ref)(long double) = ref_func; \
|
||||
UnaryDoublePrecisionBruteForceTest(kern_name##_kernel<double, lt>, ref, \
|
||||
EqValidatorBuilderFactory<lt>(), \
|
||||
static_cast<double>(std::numeric_limits<lt>::lowest()), \
|
||||
static_cast<double>(std::numeric_limits<lt>::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 <typename T>
|
||||
__global__ void remquo_kernel(std::pair<T, int>* 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<float, T>) {
|
||||
ys[i].first = remquof(x1s[i], x2s[i], &ys[i].second);
|
||||
} else if constexpr (std::is_same_v<double, T>) {
|
||||
ys[i].first = remquo(x1s[i], x2s[i], &ys[i].second);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T> std::pair<T, int> remquo_wrapper(T x1, T x2) {
|
||||
std::pair<T, int> 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<TestType>;
|
||||
std::pair<RT, int> (*ref)(RT, RT) = remquo_wrapper;
|
||||
const auto ulp_builder = ULPValidatorBuilderFactory<TestType>(0);
|
||||
const auto eq_builder = EqValidatorBuilderFactory<int>();
|
||||
|
||||
BinaryFloatingPointTest(remquo_kernel<TestType>, ref,
|
||||
PairValidatorBuilderFactory<TestType, int>(ulp_builder, eq_builder));
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_Device_remquo_remquof_Negative_RTC") { NegativeTestRTCWrapper<24>(kRemquo); }
|
||||
|
||||
template <typename T>
|
||||
__global__ void modf_kernel(std::pair<T, T>* 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].first = modff(xs[i], &ys[i].second);
|
||||
} else if constexpr (std::is_same_v<double, T>) {
|
||||
ys[i].first = modf(xs[i], &ys[i].second);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T> std::pair<T, T> modf_wrapper(T x) {
|
||||
std::pair<T, T> ret;
|
||||
ret.first = std::modf(x, &ret.second);
|
||||
return ret;
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_Device_modf_Accuracy_Positive - float") {
|
||||
UnarySinglePrecisionTest(
|
||||
modf_kernel<float>, modf_wrapper<double>,
|
||||
PairValidatorBuilderFactory<float>(ULPValidatorBuilderFactory<float>(0)));
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_Device_modf_Accuracy_Positive - double") {
|
||||
UnaryDoublePrecisionTest(
|
||||
modf_kernel<double>, modf_wrapper<long double>,
|
||||
PairValidatorBuilderFactory<double>(ULPValidatorBuilderFactory<double>(0)));
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_Device_modf_modff_Negative_RTC") { NegativeTestRTCWrapper<20>(kModf); }
|
||||
@@ -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()
|
||||
|
||||
Reference in New Issue
Block a user