Don't use accelerated vector element access for hip-clang. (#1796)

[ROCm/clr commit: 1d6a2fb064]
Этот коммит содержится в:
kpyzhov
2020-01-15 21:17:08 -05:00
коммит произвёл Rahul Garg
родитель 8bf4210095
Коммит 2d617bba3b
+22 -8
Просмотреть файл
@@ -105,14 +105,6 @@ THE SOFTWARE.
// Idea from https://t0rakka.silvrback.com/simd-scalar-accessor
Vector data;
__host__ __device__
Scalar_accessor() = default;
__host__ __device__
Scalar_accessor(T x) {
data[idx] = x;
}
__host__ __device__
operator T() const noexcept { return data[idx]; }
__host__ __device__
@@ -302,7 +294,13 @@ THE SOFTWARE.
union {
Native_vec_ data;
#if __HIP_CLANG_ONLY__
struct {
T x;
};
#else
hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
#endif
};
using value_type = T;
@@ -325,8 +323,15 @@ THE SOFTWARE.
union {
Native_vec_ data;
#if __HIP_CLANG_ONLY__
struct {
T x;
T y;
};
#else
hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
hip_impl::Scalar_accessor<T, Native_vec_, 1> y;
#endif
};
using value_type = T;
@@ -506,10 +511,19 @@ THE SOFTWARE.
union {
Native_vec_ data;
#if __HIP_CLANG_ONLY__
struct {
T x;
T y;
T z;
T w;
};
#else
hip_impl::Scalar_accessor<T, Native_vec_, 0> x;
hip_impl::Scalar_accessor<T, Native_vec_, 1> y;
hip_impl::Scalar_accessor<T, Native_vec_, 2> z;
hip_impl::Scalar_accessor<T, Native_vec_, 3> w;
#endif
};
using value_type = T;