Revert "Revert "Switch over to using native vector types, for better codegen. Remove noise.""

This reverts commit d137271083.

hipamd/include/hip/hcc_detail/texture_functions.h

@@ -110,47 +110,47 @@ union TData {
 
 #define TEXTURE_RETURN_UNSIGNED return texel.u.x;
 
-#define TEXTURE_RETURN_CHAR_X return char1(texel.i.x);
+#define TEXTURE_RETURN_CHAR_X return make_char1(texel.i.x);
 
-#define TEXTURE_RETURN_UCHAR_X return uchar1(texel.u.x);
+#define TEXTURE_RETURN_UCHAR_X return make_uchar1(texel.u.x);
 
-#define TEXTURE_RETURN_SHORT_X return short1(texel.i.x);
+#define TEXTURE_RETURN_SHORT_X return make_short1(texel.i.x);
 
-#define TEXTURE_RETURN_USHORT_X return ushort1(texel.u.x);
+#define TEXTURE_RETURN_USHORT_X return make_ushort1(texel.u.x);
 
-#define TEXTURE_RETURN_INT_X return int1(texel.i.x);
+#define TEXTURE_RETURN_INT_X return make_int1(texel.i.x);
 
-#define TEXTURE_RETURN_UINT_X return uint1(texel.u.x);
+#define TEXTURE_RETURN_UINT_X return make_uint1(texel.u.x);
 
-#define TEXTURE_RETURN_FLOAT_X return float1(texel.f.x);
+#define TEXTURE_RETURN_FLOAT_X return make_float1(texel.f.x);
 
-#define TEXTURE_RETURN_CHAR_XY return char2(texel.i.x, texel.i.y);
+#define TEXTURE_RETURN_CHAR_XY return make_char2(texel.i.x, texel.i.y);
 
-#define TEXTURE_RETURN_UCHAR_XY return uchar2(texel.u.x, texel.u.y);
+#define TEXTURE_RETURN_UCHAR_XY return make_uchar2(texel.u.x, texel.u.y);
 
-#define TEXTURE_RETURN_SHORT_XY return short2(texel.i.x, texel.i.y);
+#define TEXTURE_RETURN_SHORT_XY return make_short2(texel.i.x, texel.i.y);
 
-#define TEXTURE_RETURN_USHORT_XY return ushort2(texel.u.x, texel.u.y);
+#define TEXTURE_RETURN_USHORT_XY return make_ushort2(texel.u.x, texel.u.y);
 
-#define TEXTURE_RETURN_INT_XY return int2(texel.i.x, texel.i.y);
+#define TEXTURE_RETURN_INT_XY return make_int2(texel.i.x, texel.i.y);
 
-#define TEXTURE_RETURN_UINT_XY return uint2(texel.u.x, texel.u.y);
+#define TEXTURE_RETURN_UINT_XY return make_uint2(texel.u.x, texel.u.y);
 
-#define TEXTURE_RETURN_FLOAT_XY return float2(texel.f.x, texel.f.y);
+#define TEXTURE_RETURN_FLOAT_XY return make_float2(texel.f.x, texel.f.y);
 
-#define TEXTURE_RETURN_CHAR_XYZW return char4(texel.i.x, texel.i.y, texel.i.z, texel.i.w);
+#define TEXTURE_RETURN_CHAR_XYZW return make_char4(texel.i.x, texel.i.y, texel.i.z, texel.i.w);
 
-#define TEXTURE_RETURN_UCHAR_XYZW return uchar4(texel.u.x, texel.u.y, texel.u.z, texel.u.w);
+#define TEXTURE_RETURN_UCHAR_XYZW return make_uchar4(texel.u.x, texel.u.y, texel.u.z, texel.u.w);
 
-#define TEXTURE_RETURN_SHORT_XYZW return short4(texel.i.x, texel.i.y, texel.i.z, texel.i.w);
+#define TEXTURE_RETURN_SHORT_XYZW return make_short4(texel.i.x, texel.i.y, texel.i.z, texel.i.w);
 
-#define TEXTURE_RETURN_USHORT_XYZW return ushort4(texel.u.x, texel.u.y, texel.u.z, texel.u.w);
+#define TEXTURE_RETURN_USHORT_XYZW return make_ushort4(texel.u.x, texel.u.y, texel.u.z, texel.u.w);
 
-#define TEXTURE_RETURN_INT_XYZW return int4(texel.i.x, texel.i.y, texel.i.z, texel.i.w);
+#define TEXTURE_RETURN_INT_XYZW return make_int4(texel.i.x, texel.i.y, texel.i.z, texel.i.w);
 
-#define TEXTURE_RETURN_UINT_XYZW return uint4(texel.u.x, texel.u.y, texel.u.z, texel.u.w);
+#define TEXTURE_RETURN_UINT_XYZW return make_uint4(texel.u.x, texel.u.y, texel.u.z, texel.u.w);
 
-#define TEXTURE_RETURN_FLOAT_XYZW return float4(texel.f.x, texel.f.y, texel.f.z, texel.f.w);
+#define TEXTURE_RETURN_FLOAT_XYZW return make_float4(texel.f.x, texel.f.y, texel.f.z, texel.f.w);
 
 extern "C" {
 hc::short_vector::float4::vector_value_type __ockl_image_sample_1D(unsigned int ADDRESS_SPACE_CONSTANT* i,

hipamd/tests/src/deviceLib/vector_test_common.h

@@ -0,0 +1,105 @@
+/*
+Copyright (c) 2015-2017 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
+
+#include <type_traits>
+
+template<bool b, typename T = void>
+using Enable_if_t = typename std::enable_if<b, T>::type;
+
+__host__ __device__
+std::false_type is_vec4(...);
+__host__ __device__
+std::false_type is_vec3(...);
+__host__ __device__
+std::false_type is_vec2(...);
+__host__ __device__
+std::false_type is_vec1(...);
+
+template<typename T>
+__host__ __device__
+auto is_vec4(const T&) -> decltype(std::declval<T>().xyzw, std::true_type{});
+template<
+    typename T, Enable_if_t<decltype(!is_vec4(std::declval<T>())){}>* = nullptr>
+__host__ __device__
+auto is_vec3(const T&) -> decltype(std::declval<T>().xyz, std::true_type{});
+template<
+    typename T,
+    Enable_if_t<
+        !decltype(is_vec4(std::declval<T>())){} &&
+        !decltype(is_vec3(std::declval<T>())){}>* = nullptr>
+__host__ __device__
+auto is_vec2(const T&) -> decltype(std::declval<T>().xy, std::true_type{});
+template<
+    typename T,
+    Enable_if_t<
+        !decltype(is_vec4(std::declval<T>())){} &&
+        !decltype(is_vec3(std::declval<T>())){} &&
+        !decltype(is_vec2(std::declval<T>())){}>* = nullptr>
+__host__ __device__
+auto is_vec1(const T&) -> decltype(std::declval<T>().x, std::true_type{});
+
+template<typename T, int dimension>
+__host__ __device__
+constexpr
+bool is_vec() {
+    return (dimension == 1) ? decltype(is_vec1(std::declval<T>())){} :
+        ((dimension == 2) ? decltype(is_vec2(std::declval<T>())){} :
+        ((dimension == 3) ? decltype(is_vec3(std::declval<T>())){} :
+            decltype(is_vec4(std::declval<T>())){}));
+}
+
+template<typename T, typename U, Enable_if_t<is_vec<T, 1>()>* = nullptr>
+__host__ __device__
+inline
+bool cmp(const T& x, U expected) {
+    const auto r = x == T(expected);
+
+    return r.x != 0;
+}
+
+template<typename T, typename U, Enable_if_t<is_vec<T, 2>()>* = nullptr>
+__host__ __device__
+inline
+bool cmp(const T& x, U expected) {
+    const auto r = x == T(expected);
+
+    return r.x != 0 && r.y != 0;
+}
+
+template<typename T, typename U, Enable_if_t<is_vec<T, 3>()>* = nullptr>
+__host__ __device__
+inline
+bool cmp(const T& x, U expected) {
+    const auto r = x == T(expected);
+
+    return r.x != 0 && r.y != 0 && r.z != 0;
+}
+
+template<typename T, typename U, Enable_if_t<is_vec<T, 4>()>* = nullptr>
+__host__ __device__
+inline
+bool cmp(const T& x, U expected) {
+    const auto r = x == T(expected);
+
+    return r.x != 0 && r.y != 0 && r.z != 0 && r.w != 0;
+}