AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

EXSWHTEC-332 - Implement tests for half precision type to/from integer casting intrinsics #417

Просмотр исходного кода 
Change-Id: Icee61126274f49f3af362e35c435a9d764378f93
Этот коммит содержится в:
Nives Vukovic
2024-02-06 20:22:07 +05:30
коммит произвёл Rakesh Roy
родитель b9ea146b28
Коммит f5c3cdaaeb
5 изменённых файлов: 1017 добавлений и 0 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
catch/unit/math/CMakeLists.txt
+11
Просмотреть файл
@@ -36,6 +36,8 @@ set(TEST_SRC
half_precision_math.cc
half_precision_arithmetic.cc
half_precision_comparison.cc
casting_half2int_funcs.cc
casting_int2half_funcs.cc
)
if(HIP_PLATFORM MATCHES "nvidia")
@@ -137,3 +139,12 @@ add_test(NAME Unit_Half_Precision_Comparison_Negative
COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py
${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH}
half_precision_comparison_negative_kernels.cc 168)
add_test(NAME Unit_Device_casting_half2int_Negative
COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py
${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH}
casting_half2int_negative_kernels.cc 78)
add_test(NAME Unit_Device_casting_int2half_Negative
COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py
${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH}
casting_int2half_negative_kernels.cc 78)
catch/unit/math/casting_half2int_funcs.cc
+440
Просмотреть файл
@@ -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 "half_precision_common.hh"
#include "casting_common.hh"
/**
* @addtogroup HalfPrecisionCastingIntTypes HalfPrecisionCastingIntTypes
* @{
* @ingroup MathTest
*/
#define CAST_HALF2INT_RN_TEST_DEF(kern_name, T) \
CAST_KERNEL_DEF(kern_name, T, Float16) \
CAST_F2I_RZ_REF_DEF(kern_name, T, Float16) \
\
TEST_CASE("Unit_Device_" #kern_name "_Accuracy_Positive") { \
T (*ref)(Float16) = kern_name##_ref; \
CastUnaryHalfPrecisionTest(kern_name##_kernel, ref, EqValidatorBuilderFactory<T>()); \
}
/**
* Test Description
* ------------------------
* - Tests that checks `__half2int_rn` for all possible inputs. The results are compared against
* reference function which performs __half cast to int.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2int_rn, int)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2int_rz` for all possible inputs. The results are compared against
* reference function which performs __half cast to int.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2int_rz, int)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2int_rd` for all possible inputs. The results are compared against
* reference function which performs __half cast to int.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2int_rd, int)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2int_ru` for all possible inputs. The results are compared against
* reference function which performs __half cast to int.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2int_ru, int)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2uint_rn` for all possible inputs. The results are compared against
* reference function which performs __half cast to unsigned int.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2uint_rn, unsigned int)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2uint_rz` for all possible inputs. The results are compared against
* reference function which performs __half cast to unsigned int.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2uint_rz, unsigned int)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2uint_rd` for all possible inputs. The results are compared against
* reference function which performs __half cast to unsigned int.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2uint_rd, unsigned int)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2uint_ru` for all possible inputs. The results are compared against
* reference function which performs __half cast to unsigned int.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2uint_ru, unsigned int)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2short_rn` for all possible inputs. The results are compared
* against reference function which performs __half cast to short.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2short_rn, short)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2short_rz` for all possible inputs. The results are compared
* against reference function which performs __half cast to short.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2short_rz, short)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2short_rd` for all possible inputs. The results are compared
* against reference function which performs __half cast to short.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2short_rd, short)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2short_ru` for all possible inputs. The results are compared
* against reference function which performs __half cast to short.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2short_ru, short)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2ushort_rn` for all possible inputs. The results are compared
* against reference function which performs __half cast to unsigned short.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2ushort_rn, unsigned short)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2ushort_rz` for all possible inputs. The results are compared
* against reference function which performs __half cast to unsigned short.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2ushort_rz, unsigned short)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2ushort_rd` for all possible inputs. The results are compared
* against reference function which performs __half cast to unsigned short.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2ushort_rd, unsigned short)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2ushort_ru` for all possible inputs. The results are compared
* against reference function which performs __half cast to unsigned short.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2ushort_ru, unsigned short)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2ll_rn` for all possible inputs. The results are compared against
* reference function which performs __half cast to long long.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2ll_rn, long long)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2ll_rz` for all possible inputs. The results are compared against
* reference function which performs __half cast to long long.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2ll_rz, long long)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2ll_rd` for all possible inputs. The results are compared against
* reference function which performs __half cast to long long.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2ll_rd, long long)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2ll_ru` for all possible inputs. The results are compared against
* reference function which performs __half cast to long long.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2ll_ru, long long)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2ull_rn` for all possible inputs. The results are compared against
* reference function which performs __half cast to unsigned long long.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2ull_rn, unsigned long long)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2ull_rz` for all possible inputs. The results are compared against
* reference function which performs __half cast to unsigned long long.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2ull_rz, unsigned long long)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2ull_rd` for all possible inputs. The results are compared against
* reference function which performs __half cast to unsigned long long.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2ull_rd, unsigned long long)
/**
* Test Description
* ------------------------
* - Tests that checks `__half2ull_ru` for all possible inputs. The results are compared against
* reference function which performs __half cast to unsigned long long.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_HALF2INT_RN_TEST_DEF(__half2ull_ru, unsigned long long)
CAST_KERNEL_DEF(__half_as_short, short, Float16)
/**
* Test Description
* ------------------------
* - Tests that checks `__half_as_short` for all possible inputs. The results are compared
* against reference function which performs copy of __half value to short variable.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___half_as_short_Accuracy_Positive") {
short (*ref)(Float16) = type2_as_type1_ref<short, Float16>;
CastUnaryHalfPrecisionTest(__half_as_short_kernel, ref, EqValidatorBuilderFactory<short>());
}
CAST_KERNEL_DEF(__half_as_ushort, unsigned short, Float16)
/**
* Test Description
* ------------------------
* - Tests that checks `__half_as_ushort` for all possible inputs. The results are compared
* against reference function which performs copy of __half value to unsigned short variable.
*
* Test source
* ------------------------
* - unit/math/casting_half2int_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___half_as_ushort_Accuracy_Positive") {
unsigned short (*ref)(Float16) = type2_as_type1_ref<unsigned short, Float16>;
CastUnaryHalfPrecisionTest(__half_as_ushort_kernel, ref,
EqValidatorBuilderFactory<unsigned short>());
}
catch/unit/math/casting_half2int_negative_kernels.cc
+59
Просмотреть файл
@@ -0,0 +1,59 @@
/*
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 <hip_test_common.hh>
#include <hip/hip_fp16.h>
class Dummy {
public:
__device__ Dummy() {}
__device__ ~Dummy() {}
};
#define NEGATIVE_KERNELS_SHELL(func_name, T) \
__global__ void func_name##_kernel_v1(T* result, __half* 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, __half x) { *result = unc_name(x); }
NEGATIVE_KERNELS_SHELL(__half2int_rn, int)
NEGATIVE_KERNELS_SHELL(__half2int_rz, int)
NEGATIVE_KERNELS_SHELL(__half2int_rd, int)
NEGATIVE_KERNELS_SHELL(__half2int_ru, int)
NEGATIVE_KERNELS_SHELL(__half2uint_rn, unsigned int)
NEGATIVE_KERNELS_SHELL(__half2uint_rz, unsigned int)
NEGATIVE_KERNELS_SHELL(__half2uint_rd, unsigned int)
NEGATIVE_KERNELS_SHELL(__half2uint_ru, unsigned int)
NEGATIVE_KERNELS_SHELL(__half2short_rn, short)
NEGATIVE_KERNELS_SHELL(__half2short_rz, short)
NEGATIVE_KERNELS_SHELL(__half2short_rd, short)
NEGATIVE_KERNELS_SHELL(__half2short_ru, short)
NEGATIVE_KERNELS_SHELL(__half_as_short, short)
NEGATIVE_KERNELS_SHELL(__half2ushort_rn, unsigned short)
NEGATIVE_KERNELS_SHELL(__half2ushort_rz, unsigned short)
NEGATIVE_KERNELS_SHELL(__half2ushort_rd, unsigned short)
NEGATIVE_KERNELS_SHELL(__half2ushort_ru, unsigned short)
NEGATIVE_KERNELS_SHELL(__half_as_ushort, unsigned short)
NEGATIVE_KERNELS_SHELL(__half2ll_rn, long long)
NEGATIVE_KERNELS_SHELL(__half2ll_rz, long long)
NEGATIVE_KERNELS_SHELL(__half2ll_rd, long long)
NEGATIVE_KERNELS_SHELL(__half2ll_ru, long long)
NEGATIVE_KERNELS_SHELL(__half2ull_rn, unsigned long long)
NEGATIVE_KERNELS_SHELL(__half2ull_rz, unsigned long long)
NEGATIVE_KERNELS_SHELL(__half2ull_rd, unsigned long long)
NEGATIVE_KERNELS_SHELL(__half2ull_ru, unsigned long long)
catch/unit/math/casting_int2half_funcs.cc
+448
Просмотреть файл
@@ -0,0 +1,448 @@
/*
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 "half_precision_common.hh"
#include "casting_common.hh"
/**
* @addtogroup HalfPrecisionCastingIntTypes HalfPrecisionCastingIntTypes
* @{
* @ingroup MathTest
*/
#define CAST_INT2HALF_RN_TEST_DEF(kern_name, T) \
CAST_KERNEL_DEF(kern_name, Float16, T) \
CAST_REF_DEF(kern_name, Float16, T) \
\
TEST_CASE("Unit_Device_" #kern_name "_Accuracy_Positive") { \
Float16 (*ref)(T) = kern_name##_ref; \
CastIntRangeTest(kern_name##_kernel, ref, EqValidatorBuilderFactory<Float16>()); \
}
/**
* Test Description
* ------------------------
* - Tests that checks `__int2half_rn` for all possible inputs. The results are compared against
* reference function which performs int cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_INT2HALF_RN_TEST_DEF(__int2half_rn, int)
/**
* Test Description
* ------------------------
* - Tests that checks `__int2half_rz` for all possible inputs. The results are compared against
* reference function which performs int cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_INT2HALF_RN_TEST_DEF(__int2half_rz, int)
/**
* Test Description
* ------------------------
* - Tests that checks `__int2half_rd` for all possible inputs. The results are compared against
* reference function which performs int cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_INT2HALF_RN_TEST_DEF(__int2half_rd, int)
/**
* Test Description
* ------------------------
* - Tests that checks `__int2half_ru` for all possible inputs. The results are compared against
* reference function which performs int cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_INT2HALF_RN_TEST_DEF(__int2half_ru, int)
/**
* Test Description
* ------------------------
* - Tests that checks `__uint2half_rn` for all possible inputs. The results are compared against
* reference function which performs unsigned int cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_INT2HALF_RN_TEST_DEF(__uint2half_rn, unsigned int)
/**
* Test Description
* ------------------------
* - Tests that checks `__uint2half_rz` for all possible inputs. The results are compared against
* reference function which performs unsigned int cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_INT2HALF_RN_TEST_DEF(__uint2half_rz, unsigned int)
/**
* Test Description
* ------------------------
* - Tests that checks `__uint2half_rd` for all possible inputs. The results are compared against
* reference function which performs unsigned int cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_INT2HALF_RN_TEST_DEF(__uint2half_rd, unsigned int)
/**
* Test Description
* ------------------------
* - Tests that checks `__uint2half_ru` for all possible inputs. The results are compared against
* reference function which performs unsigned int cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_INT2HALF_RN_TEST_DEF(__uint2half_ru, unsigned int)
/**
* Test Description
* ------------------------
* - Tests that checks `__short2half_rn` for all possible inputs. The results are compared
* against reference function which performs short cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_INT2HALF_RN_TEST_DEF(__short2half_rn, short)
/**
* Test Description
* ------------------------
* - Tests that checks `__short2half_rz` for all possible inputs. The results are compared
* against reference function which performs short cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_INT2HALF_RN_TEST_DEF(__short2half_rz, short)
/**
* Test Description
* ------------------------
* - Tests that checks `__short2half_rd` for all possible inputs. The results are compared
* against reference function which performs short cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_INT2HALF_RN_TEST_DEF(__short2half_rd, short)
/**
* Test Description
* ------------------------
* - Tests that checks `__short2half_ru` for all possible inputs. The results are compared
* against reference function which performs short cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_INT2HALF_RN_TEST_DEF(__short2half_ru, short)
/**
* Test Description
* ------------------------
* - Tests that checks `__ushort2half_rn` for all possible inputs. The results are compared
* against reference function which performs unsigned short cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_INT2HALF_RN_TEST_DEF(__ushort2half_rn, unsigned short)
/**
* Test Description
* ------------------------
* - Tests that checks `__ushort2half_rz` for all possible inputs. The results are compared
* against reference function which performs unsigned short cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_INT2HALF_RN_TEST_DEF(__ushort2half_rz, unsigned short)
/**
* Test Description
* ------------------------
* - Tests that checks `__ushort2half_rd` for all possible inputs. The results are compared
* against reference function which performs unsigned short cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_INT2HALF_RN_TEST_DEF(__ushort2half_rd, unsigned short)
/**
* Test Description
* ------------------------
* - Tests that checks `__ushort2half_ru` for all possible inputs. The results are compared
* against reference function which performs unsigned short cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_INT2HALF_RN_TEST_DEF(__ushort2half_ru, unsigned short)
#define CAST_LL2HALF_TEST_DEF(kern_name, T) \
CAST_KERNEL_DEF(kern_name, Float16, T) \
CAST_REF_DEF(kern_name, Float16, T) \
\
TEST_CASE("Unit_Device_" #kern_name "_Accuracy_Positive") { \
Float16 (*ref)(T) = kern_name##_ref; \
CastIntBruteForceTest(kern_name##_kernel, ref, EqValidatorBuilderFactory<Float16>()); \
}
/**
* Test Description
* ------------------------
* - Tests that checks `__ll2half_rn` against a large number of randomly generated values. The
* results are compared against reference function which performs long long cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_LL2HALF_TEST_DEF(__ll2half_rn, long long)
/**
* Test Description
* ------------------------
* - Tests that checks `__ll2half_rz` against a large number of randomly generated values. The
* results are compared against reference function which performs long long cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_LL2HALF_TEST_DEF(__ll2half_rz, long long)
/**
* Test Description
* ------------------------
* - Tests that checks `__ll2half_rd` against a large number of randomly generated values. The
* results are compared against reference function which performs long long cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_LL2HALF_TEST_DEF(__ll2half_rd, long long)
/**
* Test Description
* ------------------------
* - Tests that checks `__ll2half_ru` against a large number of randomly generated values. The
* results are compared against reference function which performs long long cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_LL2HALF_TEST_DEF(__ll2half_ru, long long)
/**
* Test Description
* ------------------------
* - Tests that checks `__ull2half_rn` against a large number of randomly generated values. The
* results are compared against reference function which performs unsigned long long cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_LL2HALF_TEST_DEF(__ull2half_rn, unsigned long long)
/**
* Test Description
* ------------------------
* - Tests that checks `__ull2half_rz` against a large number of randomly generated values. The
* results are compared against reference function which performs unsigned long long cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_LL2HALF_TEST_DEF(__ull2half_rz, unsigned long long)
/**
* Test Description
* ------------------------
* - Tests that checks `__ull2half_rd` against a large number of randomly generated values. The
* results are compared against reference function which performs unsigned long long cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_LL2HALF_TEST_DEF(__ull2half_rd, unsigned long long)
/**
* Test Description
* ------------------------
* - Tests that checks `__ull2half_ru` against a large number of randomly generated values. The
* results are compared against reference function which performs unsigned long long cast to __half.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
CAST_LL2HALF_TEST_DEF(__ull2half_ru, unsigned long long)
CAST_KERNEL_DEF(__short_as_half, Float16, short)
/**
* Test Description
* ------------------------
* - Tests that checks `__short_as_half` for all possible inputs. The results are compared
* against reference function which performs copy of short value to __half variable.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___short_as_half_Accuracy_Positive") {
Float16 (*ref)(short) = type2_as_type1_ref<Float16, short>;
CastIntBruteForceTest(__short_as_half_kernel, ref, EqValidatorBuilderFactory<Float16>());
}
CAST_KERNEL_DEF(__ushort_as_half, Float16, unsigned short)
/**
* Test Description
* ------------------------
* - Tests that checks `__ushort_as_half` for all possible inputs. The results are compared
* against reference function which performs copy of unsigned short value to __half variable.
*
* Test source
* ------------------------
* - unit/math/casting_int2half_funcs.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___ushort_as_half_Accuracy_Positive") {
Float16 (*ref)(unsigned short) = type2_as_type1_ref<Float16, unsigned short>;
CastIntBruteForceTest(__ushort_as_half_kernel, ref, EqValidatorBuilderFactory<Float16>());
}
catch/unit/math/casting_int2half_negative_kernels.cc
+59
Просмотреть файл
@@ -0,0 +1,59 @@
/*
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 <hip_test_common.hh>
#include <hip/hip_fp16.h>
class Dummy {
public:
__device__ Dummy() {}
__device__ ~Dummy() {}
};
#define NEGATIVE_KERNELS_SHELL(func_name, T) \
__global__ void func_name##_kernel_v1(__half* result, T* x) { *result = func_name(x); } \
__global__ void func_name##_kernel_v2(__half* result, Dummy x) { *result = func_name(x); } \
__global__ void func_name##_kernel_v3(Dummy* result, T x) { *result = func_name(x); }
NEGATIVE_KERNELS_SHELL(__int2half_rn, int)
NEGATIVE_KERNELS_SHELL(__int2half_rz, int)
NEGATIVE_KERNELS_SHELL(__int2half_rd, int)
NEGATIVE_KERNELS_SHELL(__int2half_ru, int)
NEGATIVE_KERNELS_SHELL(__uint2half_rn, unsigned int)
NEGATIVE_KERNELS_SHELL(__uint2half_rz, unsigned int)
NEGATIVE_KERNELS_SHELL(__uint2half_rd, unsigned int)
NEGATIVE_KERNELS_SHELL(__uint2half_ru, unsigned int)
NEGATIVE_KERNELS_SHELL(__short2half_rn, short)
NEGATIVE_KERNELS_SHELL(__short2half_rz, short)
NEGATIVE_KERNELS_SHELL(__short2half_rd, short)
NEGATIVE_KERNELS_SHELL(__short2half_ru, short)
NEGATIVE_KERNELS_SHELL(__short_as_half, short)
NEGATIVE_KERNELS_SHELL(__ushort2half_rn, unsigned short)
NEGATIVE_KERNELS_SHELL(__ushort2half_rz, unsigned short)
NEGATIVE_KERNELS_SHELL(__ushort2half_rd, unsigned short)
NEGATIVE_KERNELS_SHELL(__ushort2half_ru, unsigned short)
NEGATIVE_KERNELS_SHELL(__ushort_as_half, unsigned short)
NEGATIVE_KERNELS_SHELL(__ll2half_rn, long long)
NEGATIVE_KERNELS_SHELL(__ll2half_rz, long long)
NEGATIVE_KERNELS_SHELL(__ll2half_rd, long long)
NEGATIVE_KERNELS_SHELL(__ll2half_ru, long long)
NEGATIVE_KERNELS_SHELL(__ull2half_rn, unsigned long long)
NEGATIVE_KERNELS_SHELL(__ull2half_rz, unsigned long long)
NEGATIVE_KERNELS_SHELL(__ull2half_rd, unsigned long long)
NEGATIVE_KERNELS_SHELL(__ull2half_ru, unsigned long long)