__half2 should walk like CUDA and talk like CUDA

Этот коммит содержится в:
Alex Voicu
2019-11-06 02:43:04 +02:00
родитель 6833886449
Коммит 55fd1363e2
+6 -12
Просмотреть файл
@@ -1126,8 +1126,7 @@ THE SOFTWARE.
{
auto r = static_cast<__half2_raw>(x).data ==
static_cast<__half2_raw>(y).data;
return __half2_raw{_Float16_2{
static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
return __half2_raw{__builtin_convertvector(-r, _Float16_2)};
}
inline
__device__
@@ -1135,8 +1134,7 @@ THE SOFTWARE.
{
auto r = static_cast<__half2_raw>(x).data !=
static_cast<__half2_raw>(y).data;
return __half2_raw{_Float16_2{
static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
return __builtin_convertvector(-r, _Float16_2);
}
inline
__device__
@@ -1144,8 +1142,7 @@ THE SOFTWARE.
{
auto r = static_cast<__half2_raw>(x).data <=
static_cast<__half2_raw>(y).data;
return __half2_raw{_Float16_2{
static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
return __builtin_convertvector(-r, _Float16_2);
}
inline
__device__
@@ -1153,8 +1150,7 @@ THE SOFTWARE.
{
auto r = static_cast<__half2_raw>(x).data >=
static_cast<__half2_raw>(y).data;
return __half2_raw{_Float16_2{
static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
return __builtin_convertvector(-r, _Float16_2);
}
inline
__device__
@@ -1162,8 +1158,7 @@ THE SOFTWARE.
{
auto r = static_cast<__half2_raw>(x).data <
static_cast<__half2_raw>(y).data;
return __half2_raw{_Float16_2{
static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
return __builtin_convertvector(-r, _Float16_2);
}
inline
__device__
@@ -1171,8 +1166,7 @@ THE SOFTWARE.
{
auto r = static_cast<__half2_raw>(x).data >
static_cast<__half2_raw>(y).data;
return __half2_raw{_Float16_2{
static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
return __builtin_convertvector(-r, _Float16_2);
}
inline
__device__