AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge pull request #1630 from ROCm-Developer-Tools/hotfix_volatile_accessors

Browse Source 
__half2 should walk like CUDA and talk like CUDA
This commit is contained in:
Rahul Garg
2019-11-06 10:25:03 -08:00
committed by GitHub
parent 85446dbd6a b9faa9f8ae
commit 2fa45c215a
1 changed files with 6 additions and 12 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/hip/hcc_detail/hip_fp16.h
+6 -12
View File
@@ -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 __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__