From 784ca6f43cdf2eb2dbdaa06af095c069c1264923 Mon Sep 17 00:00:00 2001
From: Siu Chi Chan <siuchi.chan@amd.com>
Date: Thu, 12 Mar 2020 00:00:13 -0400
Subject: [PATCH] add constexpr constructor for vector types

Change-Id: I45bb0537d6a24ee50b548c2fd8b4f20518764813
---
 include/hip/hcc_detail/hip_vector_types.h    | 311 ++++++++++++++-----
 tests/src/deviceLib/hipVectorTypes.cpp       |  45 +--
 tests/src/deviceLib/hipVectorTypesDevice.cpp | 137 ++++++--
 3 files changed, 375 insertions(+), 118 deletions(-)

diff --git a/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h
index 19259a3657..69525c5684 100644
--- a/include/hip/hcc_detail/hip_vector_types.h
+++ b/include/hip/hcc_detail/hip_vector_types.h
@@ -312,6 +312,21 @@ THE SOFTWARE.
 
         using value_type = T;
 
+        __host__ __device__
+        HIP_vector_base() = default;
+        __host__ __device__
+        explicit
+        constexpr
+        HIP_vector_base(T x) noexcept : data{x} {}
+        __host__ __device__
+        constexpr
+        HIP_vector_base(const HIP_vector_base&) = default;
+        __host__ __device__
+        constexpr
+        HIP_vector_base(HIP_vector_base&&) = default;
+        __host__ __device__
+        ~HIP_vector_base() = default;
+
         __host__ __device__
         HIP_vector_base& operator=(const HIP_vector_base& x) noexcept {
             #if __has_attribute(ext_vector_type)
@@ -347,6 +362,24 @@ THE SOFTWARE.
 
         using value_type = T;
 
+        __host__ __device__
+        HIP_vector_base() = default;
+        __host__ __device__
+        explicit
+        constexpr
+        HIP_vector_base(T x) noexcept : data{x, x} {}
+        __host__ __device__
+        constexpr
+        HIP_vector_base(T x, T y) noexcept : data{x, y} {}
+        __host__ __device__
+        constexpr
+        HIP_vector_base(const HIP_vector_base&) = default;
+        __host__ __device__
+        constexpr
+        HIP_vector_base(HIP_vector_base&&) = default;
+        __host__ __device__
+        ~HIP_vector_base() = default;
+
         __host__ __device__
         HIP_vector_base& operator=(const HIP_vector_base& x) noexcept {
             #if __has_attribute(ext_vector_type)
@@ -366,8 +399,8 @@ THE SOFTWARE.
             T d[3];
 
             __host__ __device__
-            constexpr
             Native_vec_() = default;
+
             __host__ __device__
             explicit
             constexpr
@@ -514,6 +547,29 @@ THE SOFTWARE.
         };
 
         using value_type = T;
+
+        __host__ __device__
+        HIP_vector_base() = default;
+        __host__ __device__
+        explicit
+        constexpr
+        HIP_vector_base(T x) noexcept : data{x, x, x} {}
+        __host__ __device__
+        constexpr
+        HIP_vector_base(T x, T y, T z) noexcept : data{x, y, z} {}
+        __host__ __device__
+        constexpr
+        HIP_vector_base(const HIP_vector_base&) = default;
+        __host__ __device__
+        constexpr
+        HIP_vector_base(HIP_vector_base&&) = default;
+        __host__ __device__
+        ~HIP_vector_base() = default;
+
+        __host__ __device__
+        HIP_vector_base& operator=(const HIP_vector_base&) = default;
+        __host__ __device__
+        HIP_vector_base& operator=(HIP_vector_base&&) = default;
     };
 
     template<typename T>
@@ -538,11 +594,29 @@ THE SOFTWARE.
             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            
+#endif
         };
 
         using value_type = T;
 
+        __host__ __device__
+        HIP_vector_base() = default;
+        __host__ __device__
+        explicit
+        constexpr
+        HIP_vector_base(T x) noexcept : data{x, x, x, x} {}
+        __host__ __device__
+        constexpr
+        HIP_vector_base(T x, T y, T z, T w) noexcept : data{x, y, z, w} {}
+        __host__ __device__
+        constexpr
+        HIP_vector_base(const HIP_vector_base&) = default;
+        __host__ __device__
+        constexpr
+        HIP_vector_base(HIP_vector_base&&) = default;
+        __host__ __device__
+        ~HIP_vector_base() = default;
+
         __host__ __device__
         HIP_vector_base& operator=(const HIP_vector_base& x) noexcept {
             #if __has_attribute(ext_vector_type)
@@ -563,49 +637,48 @@ THE SOFTWARE.
         using HIP_vector_base<T, rank>::data;
         using typename HIP_vector_base<T, rank>::Native_vec_;
 
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type() = default;
         template<
             typename U,
             typename std::enable_if<
                 std::is_convertible<U, T>{}>::type* = nullptr>
-        explicit inline __host__ __device__
+        __host__ __device__
+        explicit
+        constexpr
         HIP_vector_type(U x) noexcept
-        {
-            for (auto i = 0u; i != rank; ++i) data[i] = x;
-        }
+            : HIP_vector_base<T, rank>{static_cast<T>(x)}
+        {}
         template< // TODO: constrain based on type as well.
             typename... Us,
             typename std::enable_if<
                 (rank > 1) && sizeof...(Us) == rank>::type* = nullptr>
-        inline __host__ __device__
+        __host__ __device__
+        constexpr
         HIP_vector_type(Us... xs) noexcept
-        {
-            #if __has_attribute(ext_vector_type)
-                new (&data) Native_vec_{static_cast<T>(xs)...};
-            #else
-                new (&data) std::array<T, rank>{static_cast<T>(xs)...};
-            #endif
-        }
-        inline __host__ __device__
+            : HIP_vector_base<T, rank>{static_cast<T>(xs)...}
+        {}
+        __host__ __device__
+        constexpr
         HIP_vector_type(const HIP_vector_type&) = default;
-        inline __host__ __device__
+        __host__ __device__
+        constexpr
         HIP_vector_type(HIP_vector_type&&) = default;
-        inline __host__ __device__
+        __host__ __device__
         ~HIP_vector_type() = default;
 
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type& operator=(const HIP_vector_type&) = default;
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type& operator=(HIP_vector_type&&) = default;
 
         // Operators
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type& operator++() noexcept
         {
             return *this += HIP_vector_type{1};
         }
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type operator++(int) noexcept
         {
             auto tmp(*this);
@@ -613,12 +686,12 @@ THE SOFTWARE.
             return tmp;
         }
 
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type& operator--() noexcept
         {
             return *this -= HIP_vector_type{1};
         }
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type operator--(int) noexcept
         {
             auto tmp(*this);
@@ -626,7 +699,7 @@ THE SOFTWARE.
             return tmp;
         }
 
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type& operator+=(const HIP_vector_type& x) noexcept
         {
             data += x.data;
@@ -636,13 +709,13 @@ THE SOFTWARE.
             typename U,
             typename std::enable_if<
                 std::is_convertible<U, T>{}>::type* = nullptr>
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type& operator+=(U x) noexcept
         {
             return *this += HIP_vector_type{x};
         }
 
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type& operator-=(const HIP_vector_type& x) noexcept
         {
             data -= x.data;
@@ -652,13 +725,13 @@ THE SOFTWARE.
             typename U,
             typename std::enable_if<
                 std::is_convertible<U, T>{}>::type* = nullptr>
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type& operator-=(U x) noexcept
         {
             return *this -= HIP_vector_type{x};
         }
 
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type& operator*=(const HIP_vector_type& x) noexcept
         {
             data *= x.data;
@@ -668,13 +741,13 @@ THE SOFTWARE.
             typename U,
             typename std::enable_if<
                 std::is_convertible<U, T>{}>::type* = nullptr>
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type& operator*=(U x) noexcept
         {
             return *this *= HIP_vector_type{x};
         }
 
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type& operator/=(const HIP_vector_type& x) noexcept
         {
             data /= x.data;
@@ -684,7 +757,7 @@ THE SOFTWARE.
             typename U,
             typename std::enable_if<
                 std::is_convertible<U, T>{}>::type* = nullptr>
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type& operator/=(U x) noexcept
         {
             return *this /= HIP_vector_type{x};
@@ -693,7 +766,7 @@ THE SOFTWARE.
         template<
             typename U = T,
             typename std::enable_if<std::is_signed<U>{}>::type* = nullptr>
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type operator-() const noexcept
         {
             auto tmp(*this);
@@ -704,7 +777,7 @@ THE SOFTWARE.
         template<
             typename U = T,
             typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type operator~() const noexcept
         {
             HIP_vector_type r{*this};
@@ -715,7 +788,7 @@ THE SOFTWARE.
         template<
             typename U = T,
             typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type& operator%=(const HIP_vector_type& x) noexcept
         {
             data %= x.data;
@@ -725,7 +798,7 @@ THE SOFTWARE.
         template<
             typename U = T,
             typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type& operator^=(const HIP_vector_type& x) noexcept
         {
             data ^= x.data;
@@ -735,7 +808,7 @@ THE SOFTWARE.
         template<
             typename U = T,
             typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type& operator|=(const HIP_vector_type& x) noexcept
         {
             data |= x.data;
@@ -745,7 +818,7 @@ THE SOFTWARE.
         template<
             typename U = T,
             typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type& operator&=(const HIP_vector_type& x) noexcept
         {
             data &= x.data;
@@ -755,7 +828,7 @@ THE SOFTWARE.
         template<
             typename U = T,
             typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type& operator>>=(const HIP_vector_type& x) noexcept
         {
             data >>= x.data;
@@ -765,7 +838,7 @@ THE SOFTWARE.
         template<
             typename U = T,
             typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
-        inline __host__ __device__
+        __host__ __device__
         HIP_vector_type& operator<<=(const HIP_vector_type& x) noexcept
         {
             data <<= x.data;
@@ -774,21 +847,27 @@ THE SOFTWARE.
     };
 
     template<typename T, unsigned int n>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator+(
         const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
     {
         return HIP_vector_type<T, n>{x} += y;
     }
     template<typename T, unsigned int n, typename U>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator+(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
         return HIP_vector_type<T, n>{x} += HIP_vector_type<T, n>{y};
     }
     template<typename T, unsigned int n, typename U>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator+(
         U x, const HIP_vector_type<T, n>& y) noexcept
     {
@@ -796,21 +875,27 @@ THE SOFTWARE.
     }
 
     template<typename T, unsigned int n>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator-(
         const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
     {
         return HIP_vector_type<T, n>{x} -= y;
     }
     template<typename T, unsigned int n, typename U>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator-(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
         return HIP_vector_type<T, n>{x} -= HIP_vector_type<T, n>{y};
     }
     template<typename T, unsigned int n, typename U>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator-(
         U x, const HIP_vector_type<T, n>& y) noexcept
     {
@@ -818,21 +903,27 @@ THE SOFTWARE.
     }
 
     template<typename T, unsigned int n>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator*(
         const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
     {
         return HIP_vector_type<T, n>{x} *= y;
     }
     template<typename T, unsigned int n, typename U>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator*(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
         return HIP_vector_type<T, n>{x} *= HIP_vector_type<T, n>{y};
     }
     template<typename T, unsigned int n, typename U>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator*(
         U x, const HIP_vector_type<T, n>& y) noexcept
     {
@@ -840,64 +931,90 @@ THE SOFTWARE.
     }
 
     template<typename T, unsigned int n>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator/(
         const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
     {
         return HIP_vector_type<T, n>{x} /= y;
     }
     template<typename T, unsigned int n, typename U>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator/(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
         return HIP_vector_type<T, n>{x} /= HIP_vector_type<T, n>{y};
     }
     template<typename T, unsigned int n, typename U>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator/(
         U x, const HIP_vector_type<T, n>& y) noexcept
     {
         return HIP_vector_type<T, n>{x} /= y;
     }
 
+    template<typename V>
+    __host__ __device__
+    inline
+    constexpr
+    bool _hip_any_zero(const V& x, int n) noexcept
+    {
+        return
+            (n == -1) ? true : ((x[n] == 0) ? false : _hip_any_zero(x, n - 1));
+    }
+
     template<typename T, unsigned int n>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     bool operator==(
         const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
     {
-        auto tmp = x.data == y.data;
-        for (auto i = 0u; i != n; ++i) if (tmp[i] == 0) return false;
-        return true;
+        return _hip_any_zero(x.data == y.data, n - 1);
     }
     template<typename T, unsigned int n, typename U>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     bool operator==(const HIP_vector_type<T, n>& x, U y) noexcept
     {
         return x == HIP_vector_type<T, n>{y};
     }
     template<typename T, unsigned int n, typename U>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     bool operator==(U x, const HIP_vector_type<T, n>& y) noexcept
     {
         return HIP_vector_type<T, n>{x} == y;
     }
 
     template<typename T, unsigned int n>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     bool operator!=(
         const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
     {
         return !(x == y);
     }
     template<typename T, unsigned int n, typename U>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     bool operator!=(const HIP_vector_type<T, n>& x, U y) noexcept
     {
         return !(x == y);
     }
     template<typename T, unsigned int n, typename U>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     bool operator!=(U x, const HIP_vector_type<T, n>& y) noexcept
     {
         return !(x == y);
@@ -907,7 +1024,9 @@ THE SOFTWARE.
         typename T,
         unsigned int n,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator%(
         const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
     {
@@ -918,7 +1037,9 @@ THE SOFTWARE.
         unsigned int n,
         typename U,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator%(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
@@ -929,7 +1050,9 @@ THE SOFTWARE.
         unsigned int n,
         typename U,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator%(
         U x, const HIP_vector_type<T, n>& y) noexcept
     {
@@ -940,7 +1063,9 @@ THE SOFTWARE.
         typename T,
         unsigned int n,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator^(
         const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
     {
@@ -951,7 +1076,9 @@ THE SOFTWARE.
         unsigned int n,
         typename U,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator^(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
@@ -962,7 +1089,9 @@ THE SOFTWARE.
         unsigned int n,
         typename U,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator^(
         U x, const HIP_vector_type<T, n>& y) noexcept
     {
@@ -973,7 +1102,9 @@ THE SOFTWARE.
         typename T,
         unsigned int n,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator|(
         const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
     {
@@ -984,7 +1115,9 @@ THE SOFTWARE.
         unsigned int n,
         typename U,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator|(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
@@ -995,7 +1128,9 @@ THE SOFTWARE.
         unsigned int n,
         typename U,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator|(
         U x, const HIP_vector_type<T, n>& y) noexcept
     {
@@ -1006,7 +1141,9 @@ THE SOFTWARE.
         typename T,
         unsigned int n,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator&(
         const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
     {
@@ -1017,7 +1154,9 @@ THE SOFTWARE.
         unsigned int n,
         typename U,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator&(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
@@ -1028,7 +1167,9 @@ THE SOFTWARE.
         unsigned int n,
         typename U,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator&(
         U x, const HIP_vector_type<T, n>& y) noexcept
     {
@@ -1039,7 +1180,9 @@ THE SOFTWARE.
         typename T,
         unsigned int n,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator>>(
         const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
     {
@@ -1050,7 +1193,9 @@ THE SOFTWARE.
         unsigned int n,
         typename U,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator>>(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
@@ -1061,7 +1206,9 @@ THE SOFTWARE.
         unsigned int n,
         typename U,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator>>(
         U x, const HIP_vector_type<T, n>& y) noexcept
     {
@@ -1072,7 +1219,9 @@ THE SOFTWARE.
         typename T,
         unsigned int n,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator<<(
         const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
     {
@@ -1083,7 +1232,9 @@ THE SOFTWARE.
         unsigned int n,
         typename U,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator<<(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
@@ -1095,7 +1246,9 @@ THE SOFTWARE.
         typename U,
         typename std::enable_if<std::is_arithmetic<U>::value>::type,
         typename std::enable_if<std::is_integral<T>{}>* = nullptr>
-    inline __host__ __device__
+    __host__ __device__
+    inline
+    constexpr
     HIP_vector_type<T, n> operator<<(
         U x, const HIP_vector_type<T, n>& y) noexcept
     {
diff --git a/tests/src/deviceLib/hipVectorTypes.cpp b/tests/src/deviceLib/hipVectorTypes.cpp
index 70c8320073..31bd7ee08c 100644
--- a/tests/src/deviceLib/hipVectorTypes.cpp
+++ b/tests/src/deviceLib/hipVectorTypes.cpp
@@ -112,49 +112,54 @@ bool constructor_tests() {
 
 template<typename V>
 bool TestVectorType() {
+    constexpr V v1{1};
+    constexpr V v2{2};
+    constexpr V v3{3};
+    constexpr V v4{4};
+
     V f1{1};
     V f2{1};
     V f3 = f1 + f2;
-    if (f3 != V{2}) return false;
+    if (f3 != v2) return false;
     f2 = f3 - f1;
-    if (f2 != V{1}) return false;
+    if (f2 != v1) return false;
     f1 = f2 * f3;
-    if (f1 != V{2}) return false;
+    if (f1 != v2) return false;
     f2 = f1 / f3;
-    if (f2 != V{1}) return false;
+    if (f2 != v1) return false;
     if (!integer_binary_tests(f1, f2, f3)) return false;
 
     f1 = V{2};
     f2 = V{1};
     f1 += f2;
-    if (f1 != V{3}) return false;
+    if (f1 != v3) return false;
     f1 -= f2;
-    if (f1 != V{2}) return false;
+    if (f1 != v2) return false;
     f1 *= f2;
-    if (f1 != V{2}) return false;
+    if (f1 != v2) return false;
     f1 /= f2;
-    if (f1 != V{2}) return false;
+    if (f1 != v2) return false;
     if (!integer_unary_tests(f1, f2)) return false;
 
-    f1 = V{2};
+    f1 = v2;
     f2 = f1++;
-    if (f1 != V{3}) return false;
-    if (f2 != V{2}) return false;
+    if (f1 != v3) return false;
+    if (f2 != v2) return false;
     f2 = f1--;
-    if (f2 != V{3}) return false;
-    if (f1 != V{2}) return false;
+    if (f2 != v3) return false;
+    if (f1 != v2) return false;
     f2 = ++f1;
-    if (f1 != V{3}) return false;
-    if (f2 != V{3}) return false;
+    if (f1 != v3) return false;
+    if (f2 != v3) return false;
     f2 = --f1;
-    if (f1 != V{2}) return false;
-    if (f2 != V{2}) return false;
+    if (f1 != v2) return false;
+    if (f2 != v2) return false;
 
     if (!constructor_tests<V>()) return false;
 
-    f1 = V{3};
-    f2 = V{4};
-    f3 = V{3};
+    f1 = v3;
+    f2 = v4;
+    f3 = v3;
     if (f1 == f2) return false;
     if (!(f1 != f2)) return false;
 
diff --git a/tests/src/deviceLib/hipVectorTypesDevice.cpp b/tests/src/deviceLib/hipVectorTypesDevice.cpp
index 4bf5d2c87d..ec3baa73ad 100644
--- a/tests/src/deviceLib/hipVectorTypesDevice.cpp
+++ b/tests/src/deviceLib/hipVectorTypesDevice.cpp
@@ -105,6 +105,11 @@ bool integer_binary_tests(V& f1, V& f2, V& f3) {
 template<typename V>
 __device__
 bool TestVectorType() {
+    constexpr V v1{1};
+    constexpr V v2{2};
+    constexpr V v3{3};
+    constexpr V v4{4};
+
     V f1{1};
     V f2{1};
     V f3 = f1 + f2;
@@ -117,41 +122,41 @@ bool TestVectorType() {
     if (f2 != V{1}) return false;
     if (!integer_binary_tests(f1, f2, f3)) return false;
 
-    f1 = V{2};
-    f2 = V{1};
+    f1 = v2;
+    f2 = v1;
     f1 += f2;
-    if (f1 != V{3}) return false;
+    if (f1 != v3) return false;
     f1 -= f2;
-    if (f1 != V{2}) return false;
+    if (f1 != v2) return false;
     f1 *= f2;
-    if (f1 != V{2}) return false;
+    if (f1 != v2) return false;
     f1 /= f2;
-    if (f1 != V{2}) return false;
+    if (f1 != v2) return false;
     if (!integer_unary_tests(f1, f2)) return false;
 
-    f1 = V{2};
+    f1 = v2;
     f2 = f1++;
-    if (f1 != V{3}) return false;
-    if (f2 != V{2}) return false;
+    if (f1 != v3) return false;
+    if (f2 != v2) return false;
     f2 = f1--;
-    if (f2 != V{3}) return false;
-    if (f1 != V{2}) return false;
+    if (f2 != v3) return false;
+    if (f1 != v2) return false;
     f2 = ++f1;
-    if (f1 != V{3}) return false;
-    if (f2 != V{3}) return false;
+    if (f1 != v3) return false;
+    if (f2 != v3) return false;
     f2 = --f1;
-    if (f1 != V{2}) return false;
-    if (f2 != V{2}) return false;
+    if (f1 != v2) return false;
+    if (f2 != v2) return false;
 
-    f1 = V{3};
-    f2 = V{4};
-    f3 = V{3};
+    f1 = v3;
+    f2 = v4;
+    f3 = v3;
     if (f1 == f2) return false;
     if (!(f1 != f2)) return false;
 
     #if 0 // TODO: investigate on GFX8
         using T = typename V::value_type;
-    
+
         const T& x = f1.x;
         T& y = f2.x;
         const volatile T& z = f3.x;
@@ -196,6 +201,86 @@ void CheckVectorTypes(bool* ptr) {
         double1, double2, double3, double4>();
 }
 
+
+template<typename V>
+__global__
+void CheckSharedVectorType(bool* ptr) {
+    constexpr V v1{1};
+    constexpr V v2{2};
+    constexpr V v3{3};
+    constexpr V v4{4};
+    __shared__ V f1, f2, f3;
+
+    *ptr = true;
+    f1 = V{1};
+    f2 = V{1};
+    f3 = f1 + f2;
+    *ptr = *ptr && f3 == V{2};
+    f2 = f3 - f1;
+    *ptr = *ptr && f2 == V{1};
+    f1 = f2 * f3;
+    *ptr = *ptr && f1 == V{2};
+    f2 = f1 / f3;
+    *ptr = *ptr && f2 == V{1};
+    *ptr = *ptr && integer_binary_tests(f1, f2, f3);
+
+    f1 = v2;
+    f2 = v1;
+    f1 += f2;
+    *ptr = *ptr && f1 == v3;
+    f1 -= f2;
+    *ptr = *ptr && f1 == v2;
+    f1 *= f2;
+    *ptr = *ptr && f1 == v2;
+    f1 /= f2;
+    *ptr = *ptr && f1 == v2;
+    *ptr = *ptr && integer_unary_tests(f1, f2);
+
+    f1 = v2;
+    f2 = f1++;
+    *ptr = *ptr && f1 == v3;
+    *ptr = *ptr && f2 == v2;
+    f2 = f1--;
+    *ptr = *ptr && f2 == v3;
+    *ptr = *ptr && f1 == v2;
+    f2 = ++f1;
+    *ptr = *ptr && f1 == v3;
+    *ptr = *ptr && f2 == v3;
+    f2 = --f1;
+    *ptr = *ptr && f1 == v2;
+    *ptr = *ptr && f2 == v2;
+
+    f1 = v3;
+    f2 = v4;
+    f3 = v3;
+    *ptr = *ptr && f1 != f2;
+}
+
+template <typename V>
+bool run_CheckSharedVectorType() {
+    bool* ptr = nullptr;
+    if (hipMalloc(&ptr, sizeof(bool)) != HIP_SUCCESS) return false;
+    unique_ptr<bool, decltype(hipFree)*> correct{ptr, hipFree};
+    hipLaunchKernelGGL(
+        (CheckSharedVectorType<V>), dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, correct.get());
+    bool passed = true;
+    if (hipMemcpyDtoH(&passed, correct.get(), sizeof(bool)) != HIP_SUCCESS) {
+        return false;
+    }
+    return passed;
+}
+
+template<typename... Ts, Enable_if_t<sizeof...(Ts) == 0>* = nullptr>
+bool run_CheckSharedVectorTypes() {
+    return true;
+}
+
+template <typename V, typename... Vs>
+bool run_CheckSharedVectorTypes() {
+    return run_CheckSharedVectorType<V>() &&
+             run_CheckSharedVectorTypes<Vs...>();
+}
+
 int main() {
     static_assert(sizeof(float1) == 4, "");
     static_assert(sizeof(float2) >= 8, "");
@@ -212,6 +297,20 @@ int main() {
         return EXIT_FAILURE;
     }
 
+    passed = passed && run_CheckSharedVectorTypes<
+        char1, char2, char3, char4,
+        uchar1, uchar2, uchar3, uchar4,
+        short1, short2, short3, short4,
+        ushort1, ushort2, ushort3, ushort4,
+        int1, int2, int3, int4,
+        uint1, uint2, uint3, uint4,
+        long1, long2, long3, long4,
+        ulong1, ulong2, ulong3, ulong4,
+        longlong1, longlong2, longlong3, longlong4,
+        ulonglong1, ulonglong2, ulonglong3, ulonglong4,
+        float1, float2, float3, float4,
+        double1, double2, double3, double4>();
+
     if (passed == true) {
         passed();
     }