diff --git a/catch/unit/math/CMakeLists.txt b/catch/unit/math/CMakeLists.txt index 347d4e10ef..b3e5a937d1 100644 --- a/catch/unit/math/CMakeLists.txt +++ b/catch/unit/math/CMakeLists.txt @@ -30,6 +30,7 @@ set(TEST_SRC log_funcs.cc special_funcs.cc casting_double_funcs.cc + casting_float_funcs.cc ) if(HIP_PLATFORM MATCHES "nvidia") @@ -106,3 +107,7 @@ add_test(NAME Unit_Device_casting_double_Negative COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} casting_double_negative_kernels.cc 69) +add_test(NAME Unit_Device_casting_float_Negative + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + casting_float_negative_kernels.cc 54) diff --git a/catch/unit/math/casting_float_funcs.cc b/catch/unit/math/casting_float_funcs.cc new file mode 100644 index 0000000000..f5e92e218b --- /dev/null +++ b/catch/unit/math/casting_float_funcs.cc @@ -0,0 +1,440 @@ +/* +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 "casting_common.hh" +#include "casting_float_negative_kernels_rtc.hh" + +/** + * @addtogroup CastingFloatType CastingFloatType + * @{ + * @ingroup MathTest + */ + +#define CAST_FLOAT2INT_TEST_DEF(kern_name, T, ref_func) \ + CAST_KERNEL_DEF(kern_name, T, float) \ + CAST_F2I_REF_DEF(kern_name, T, float, ref_func) \ + \ + TEST_CASE("Unit_Device_" #kern_name "_Positive") { \ + T (*ref)(float) = kern_name##_ref; \ + UnarySinglePrecisionRangeTest(kern_name##_kernel, ref, EqValidatorBuilderFactory(), \ + std::numeric_limits::lowest(), \ + std::numeric_limits::max()); \ + } + +#define CAST_FLOAT2INT_RZ_TEST_DEF(kern_name, T) \ + CAST_KERNEL_DEF(kern_name, T, float) \ + CAST_F2I_RZ_REF_DEF(kern_name, T, float) \ + \ + TEST_CASE("Unit_Device_" #kern_name "_Positive") { \ + T (*ref)(float) = kern_name##_ref; \ + UnarySinglePrecisionRangeTest(kern_name##_kernel, ref, EqValidatorBuilderFactory(), \ + std::numeric_limits::lowest(), \ + std::numeric_limits::max()); \ + } + +/** + * Test Description + * ------------------------ + * - Tests that checks `__float2int_rd` for all possible inputs. The results are compared against + * reference function `std::floor`. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_FLOAT2INT_TEST_DEF(__float2int_rd, int, std::floor) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__float2int_rn` for all possible inputs. The results are compared against + * reference function `std::rint`. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_FLOAT2INT_TEST_DEF(__float2int_rn, int, std::rint) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__float2int_ru` for all possible inputs. The results are compared against + * reference function `std::ceil`. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_FLOAT2INT_TEST_DEF(__float2int_ru, int, std::ceil) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__float2int_rz` for all possible inputs. The results are compared against + * reference function `std::trunc`. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_FLOAT2INT_TEST_DEF(__float2int_rz, int, std::trunc) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for __float2int_[rd,rn,ru,rz]. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___float2int_Negative_RTC") { NegativeTestRTCWrapper<12>(kFloat2Int); } + +/** + * Test Description + * ------------------------ + * - Tests that checks `__float2uint_rd` for all possible inputs. The results are compared + * against reference function `std::floor`. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_FLOAT2INT_TEST_DEF(__float2uint_rd, unsigned int, std::floor) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__float2uint_rn` for all possible inputs. The results are compared + * against reference function `std::rint`. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_FLOAT2INT_TEST_DEF(__float2uint_rn, unsigned int, std::rint) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__float2uint_ru` for all possible inputs. The results are compared + * against reference function `std::ceil`. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_FLOAT2INT_TEST_DEF(__float2uint_ru, unsigned int, std::ceil) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__float2uint_rz` against a table of difficult values, followed by a large + * number of randomly generated values. The results are compared against reference function which + * performs cast to unsigned int. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_FLOAT2INT_RZ_TEST_DEF(__float2uint_rz, unsigned int) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for __float2uint_[rd,rn,ru,rz]. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___float2uint_Negative_RTC") { NegativeTestRTCWrapper<12>(kFloat2Uint); } + +#define CAST_FLOAT2LL_TEST_DEF(kern_name, T, ref_func) \ + CAST_KERNEL_DEF(kern_name, T, float) \ + CAST_F2I_REF_DEF(kern_name, T, float, ref_func) \ + \ + TEST_CASE("Unit_Device_" #kern_name "_Positive") { \ + T (*ref)(float) = kern_name##_ref; \ + UnarySinglePrecisionRangeTest(kern_name##_kernel, ref, EqValidatorBuilderFactory(), \ + static_cast(std::numeric_limits::min()), \ + static_cast(std::numeric_limits::max())); \ + } + +#define CAST_FLOAT2LL_RZ_TEST_DEF(kern_name, T) \ + CAST_KERNEL_DEF(kern_name, T, float) \ + CAST_F2I_RZ_REF_DEF(kern_name, T, float) \ + \ + TEST_CASE("Unit_Device_" #kern_name "_Positive") { \ + T (*ref)(float) = kern_name##_ref; \ + UnarySinglePrecisionRangeTest(kern_name##_kernel, ref, EqValidatorBuilderFactory(), \ + static_cast(std::numeric_limits::min()), \ + static_cast(std::numeric_limits::max())); \ + } + +/** + * Test Description + * ------------------------ + * - Tests that checks `__float2ll_rd` for all possible inputs. The results are compared against + * reference function `std::floor`. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_FLOAT2LL_TEST_DEF(__float2ll_rd, long long int, std::floor) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__float2ll_rn` for all possible inputs between lowest and maximal long + * long int value. The results are compared against reference function `std::rint`. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_FLOAT2LL_TEST_DEF(__float2ll_rn, long long int, std::rint) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__float2ll_ru` for all possible inputs between lowest and maximal long + * long int value. The results are compared against reference function `std::ceil`. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_FLOAT2LL_TEST_DEF(__float2ll_ru, long long int, std::ceil) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__float2ll_rz` for all possible inputs between lowest and maximal long + * long int value. The results are compared against reference function which performs cast to long + * long int. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_FLOAT2LL_RZ_TEST_DEF(__float2ll_rz, long long int) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for __float2ll_[rd,rn,ru,rz]. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___float2ll_Negative_RTC") { NegativeTestRTCWrapper<12>(kFloat2LL); } + +/** + * Test Description + * ------------------------ + * - Tests that checks `__float2ull_rd` for all possible inputs between lowest and maximal + * unsigned long long int value. The results are compared against reference function `std::floor`. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_FLOAT2LL_TEST_DEF(__float2ull_rd, unsigned long long int, std::floor) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__float2ull_rn` for all possible inputs between lowest and maximal + * unsigned long long int value. The results are compared against reference function `std::rint`. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_FLOAT2LL_TEST_DEF(__float2ull_rn, unsigned long long int, std::rint) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__float2ull_ru` for all possible inputs between lowest and maximal + * unsigned long long int value. The results are compared against reference function `std::ceil`. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_FLOAT2LL_TEST_DEF(__float2ull_ru, unsigned long long int, std::ceil) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__float2ll_rz` for all possible inputs between lowest and maximal + * unsigned long long int value. The results are compared against reference function which performs + * cast to unsigned long long int. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +CAST_FLOAT2LL_RZ_TEST_DEF(__float2ull_rz, unsigned long long int) + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for __float2ull_[rd,rn,ru,rz]. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___float2ull_Negative_RTC") { NegativeTestRTCWrapper<12>(kFloat2ULL); } + +CAST_KERNEL_DEF(__float_as_int, int, float) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__float_as_int` for all possible inputs. The results are compared against + * reference function which performs copy of float value to int variable. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___float_as_int_Positive") { + int (*ref)(float) = type2_as_type1_ref; + UnarySinglePrecisionTest(__float_as_int_kernel, ref, EqValidatorBuilderFactory()); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for __float_as_int. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___float_as_int_Negative_RTC") { NegativeTestRTCWrapper<3>(kFloatAsInt); } + +CAST_KERNEL_DEF(__float_as_uint, unsigned int, float) + +/** + * Test Description + * ------------------------ + * - Tests that checks `__float_as_uint` for all possible inputs. The results are compared + * against reference function which performs copy of float value to unsigned int variable. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___float_as_uint_Positive") { + unsigned int (*ref)(float) = type2_as_type1_ref; + UnarySinglePrecisionTest(__float_as_uint_kernel, ref, EqValidatorBuilderFactory()); +} + +/** + * Test Description + * ------------------------ + * - RTCs kernels that pass argument of invalid type for __float_as_uint. + * + * Test source + * ------------------------ + * - unit/math/casting_float_funcs.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Device___float_as_uint_Negative_RTC") { NegativeTestRTCWrapper<3>(kFloatAsUint); } diff --git a/catch/unit/math/casting_float_negative_kernels.cc b/catch/unit/math/casting_float_negative_kernels.cc new file mode 100644 index 0000000000..eecbd6dd7e --- /dev/null +++ b/catch/unit/math/casting_float_negative_kernels.cc @@ -0,0 +1,50 @@ +/* +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, T) \ + __global__ void func_name##_kernel_v1(T* result, float* x) { *result = func_name(x); } \ + __global__ void func_name##_kernel_v2(T* result, Dummy x) { *result = func_name(x); } \ + __global__ void func_name##_kernel_v3(Dummy* result, float x) { *result = func_name(x); } + +NEGATIVE_KERNELS_SHELL(__float2int_rd, int) +NEGATIVE_KERNELS_SHELL(__float2int_rn, int) +NEGATIVE_KERNELS_SHELL(__float2int_ru, int) +NEGATIVE_KERNELS_SHELL(__float2int_rz, int) +NEGATIVE_KERNELS_SHELL(__float2uint_rd, unsigned int) +NEGATIVE_KERNELS_SHELL(__float2uint_rn, unsigned int) +NEGATIVE_KERNELS_SHELL(__float2uint_ru, unsigned int) +NEGATIVE_KERNELS_SHELL(__float2uint_rz, unsigned int) +NEGATIVE_KERNELS_SHELL(__float2ll_rd, long long int) +NEGATIVE_KERNELS_SHELL(__float2ll_rn, long long int) +NEGATIVE_KERNELS_SHELL(__float2ll_ru, long long int) +NEGATIVE_KERNELS_SHELL(__float2ll_rz, long long int) +NEGATIVE_KERNELS_SHELL(__float2ull_rd, unsigned long long int) +NEGATIVE_KERNELS_SHELL(__float2ull_rn, unsigned long long int) +NEGATIVE_KERNELS_SHELL(__float2ull_ru, unsigned long long int) +NEGATIVE_KERNELS_SHELL(__float2ull_rz, unsigned long long int) +NEGATIVE_KERNELS_SHELL(__float_as_int, int) +NEGATIVE_KERNELS_SHELL(__float_as_uint, unsigned int) \ No newline at end of file diff --git a/catch/unit/math/casting_float_negative_kernels_rtc.hh b/catch/unit/math/casting_float_negative_kernels_rtc.hh new file mode 100644 index 0000000000..45fba3b0e7 --- /dev/null +++ b/catch/unit/math/casting_float_negative_kernels_rtc.hh @@ -0,0 +1,126 @@ +/* +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 float type casting negative Test Cases that are using RTC. +*/ + +static constexpr auto kFloat2Int{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void float2int_rd_kernel_v1(int* result, float* x) { *result = __float2int_rd(x); } + __global__ void float2int_rd_kernel_v2(int* result, Dummy x) { *result = __float2int_rd(x); } + __global__ void float2int_rd_kernel_v3(Dummy* result, float x) { *result = __float2int_rd(x); } + __global__ void float2int_rn_kernel_v1(int* result, float* x) { *result = __float2int_rn(x); } + __global__ void float2int_rn_kernel_v2(int* result, Dummy x) { *result = __float2int_rn(x); } + __global__ void float2int_rn_kernel_v3(Dummy* result, float x) { *result = __float2int_rn(x); } + __global__ void float2int_ru_kernel_v1(int* result, float* x) { *result = __float2int_ru(x); } + __global__ void float2int_ru_kernel_v2(int* result, Dummy x) { *result = __float2int_ru(x); } + __global__ void float2int_ru_kernel_v3(Dummy* result, float x) { *result = __float2int_ru(x); } + __global__ void float2int_rz_kernel_v1(int* result, float* x) { *result = __float2int_rz(x); } + __global__ void float2int_rz_kernel_v2(int* result, Dummy x) { *result = __float2int_rz(x); } + __global__ void float2int_rz_kernel_v3(Dummy* result, float x) { *result = __float2int_rz(x); } +)"}; + +static constexpr auto kFloat2Uint{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void float2uint_rd_kernel_v1(unsigned int* result, float* x) { *result = __float2uint_rd(x); } + __global__ void float2uint_rd_kernel_v2(unsigned int* result, Dummy x) { *result = __float2uint_rd(x); } + __global__ void float2uint_rd_kernel_v3(Dummy* result, float x) { *result = __float2uint_rd(x); } + __global__ void float2uint_rn_kernel_v1(unsigned int* result, float* x) { *result = __float2uint_rn(x); } + __global__ void float2uint_rn_kernel_v2(unsigned int* result, Dummy x) { *result = __float2uint_rn(x); } + __global__ void float2uint_rn_kernel_v3(Dummy* result, float x) { *result = __float2uint_rn(x); } + __global__ void float2uint_ru_kernel_v1(unsigned int* result, float* x) { *result = __float2uint_ru(x); } + __global__ void float2uint_ru_kernel_v2(unsigned int* result, Dummy x) { *result = __float2uint_ru(x); } + __global__ void float2uint_ru_kernel_v3(Dummy* result, float x) { *result = __float2uint_ru(x); } + __global__ void float2uint_rz_kernel_v1(unsigned int* result, float* x) { *result = __float2uint_rz(x); } + __global__ void float2uint_rz_kernel_v2(unsigned int* result, Dummy x) { *result = __float2uint_rz(x); } + __global__ void float2uint_rz_kernel_v3(Dummy* result, float x) { *result = __float2uint_rz(x); } +)"}; + +static constexpr auto kFloat2LL{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void float2ll_rd_kernel_v1(long long int* result, float* x) { *result = __float2ll_rd(x); } + __global__ void float2ll_rd_kernel_v2(long long int* result, Dummy x) { *result = __float2ll_rd(x); } + __global__ void float2ll_rd_kernel_v3(Dummy* result, float x) { *result = __float2ll_rd(x); } + __global__ void float2ll_rn_kernel_v1(long long int* result, float* x) { *result = __float2ll_rn(x); } + __global__ void float2ll_rn_kernel_v2(long long int* result, Dummy x) { *result = __float2ll_rn(x); } + __global__ void float2ll_rn_kernel_v3(Dummy* result, float x) { *result = __float2ll_rn(x); } + __global__ void float2ll_ru_kernel_v1(long long int* result, float* x) { *result = __float2ll_ru(x); } + __global__ void float2ll_ru_kernel_v2(long long int* result, Dummy x) { *result = __float2ll_ru(x); } + __global__ void float2ll_ru_kernel_v3(Dummy* result, float x) { *result = __float2ll_ru(x); } + __global__ void float2ll_rz_kernel_v1(long long int* result, float* x) { *result = __float2ll_rz(x); } + __global__ void float2ll_rz_kernel_v2(long long int* result, Dummy x) { *result = __float2ll_rz(x); } + __global__ void float2ll_rz_kernel_v3(Dummy* result, float x) { *result = __float2ll_rz(x); } +)"}; + +static constexpr auto kFloat2ULL{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void float2ull_rd_kernel_v1(unsigned long long int* result, float* x) { *result = __float2ull_rd(x); } + __global__ void float2ull_rd_kernel_v2(unsigned long long int* result, Dummy x) { *result = __float2ull_rd(x); } + __global__ void float2ull_rd_kernel_v3(Dummy* result, float x) { *result = __float2ull_rd(x); } + __global__ void float2ull_rn_kernel_v1(unsigned long long int* result, float* x) { *result = __float2ull_rn(x); } + __global__ void float2ull_rn_kernel_v2(unsigned long long int* result, Dummy x) { *result = __float2ull_rn(x); } + __global__ void float2ull_rn_kernel_v3(Dummy* result, float x) { *result = __float2ull_rn(x); } + __global__ void float2ull_ru_kernel_v1(unsigned long long int* result, float* x) { *result = __float2ull_ru(x); } + __global__ void float2ull_ru_kernel_v2(unsigned long long int* result, Dummy x) { *result = __float2ull_ru(x); } + __global__ void float2ull_ru_kernel_v3(Dummy* result, float x) { *result = __float2ull_ru(x); } + __global__ void float2ull_rz_kernel_v1(unsigned long long int* result, float* x) { *result = __float2ull_rz(x); } + __global__ void float2ull_rz_kernel_v2(unsigned long long int* result, Dummy x) { *result = __float2ull_rz(x); } + __global__ void float2ull_rz_kernel_v3(Dummy* result, float x) { *result = __float2ull_rz(x); } +)"}; + +static constexpr auto kFloatAsInt{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void float_as_int_kernel_v1(int* result, float* x) { *result = __float_as_int(x); } + __global__ void float_as_int_kernel_v2(int* result, Dummy x) { *result = __float_as_int(x); } + __global__ void float_as_int_kernel_v3(Dummy* result, float x) { *result = __float_as_int(x); } +)"}; + +static constexpr auto kFloatAsUint{R"( + class Dummy { + public: + __device__ Dummy() {} + __device__ ~Dummy() {} + }; + __global__ void float_as_uint_kernel_v1(unsigned int* result, float* x) { *result = __float_as_uint(x); } + __global__ void float_as_uint_kernel_v2(unsigned int* result, Dummy x) { *result = __float_as_uint(x); } + __global__ void float_as_uint_kernel_v3(Dummy* result, float x) { *result = __float_as_uint(x); } +)"}; diff --git a/catch/unit/math/math_common.hh b/catch/unit/math/math_common.hh index 010780474f..738dbf66c0 100644 --- a/catch/unit/math/math_common.hh +++ b/catch/unit/math/math_common.hh @@ -215,7 +215,7 @@ 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"); + std::string args = std::string("-ferror-limit=200"); const char* options[] = {args.c_str()}; hiprtcResult result{hiprtcCompileProgram(program, 1, options)}; #else