From 91cc897d80b459dd4f75ebd72e962449401eaf7b Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu@amd.com>
Date: Tue, 17 Jul 2018 17:17:48 +0100
Subject: [PATCH] Constrain variadic constructor to ranks > 1 and add missing
 `__device__`s

[ROCm/hip commit: 96aa6e22ee2197ecdad16c882dabe1e0298e1dc4]
---
 projects/hip/include/hip/hcc_detail/hip_vector_types.h    | 3 ++-
 projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp | 2 ++
 2 files changed, 4 insertions(+), 1 deletion(-)

diff --git a/projects/hip/include/hip/hcc_detail/hip_vector_types.h b/projects/hip/include/hip/hcc_detail/hip_vector_types.h
index cf7058af2b..371dfad189 100644
--- a/projects/hip/include/hip/hcc_detail/hip_vector_types.h
+++ b/projects/hip/include/hip/hcc_detail/hip_vector_types.h
@@ -120,7 +120,8 @@ THE SOFTWARE.
         }
         template< // TODO: constrain based on type as well.
             typename... Us,
-            typename std::enable_if<sizeof...(Us) == rank>::type* = nullptr>
+            typename std::enable_if<
+                (rank > 1) && sizeof...(Us) == rank>::type* = nullptr>
         __host__ __device__
         HIP_vector_type(Us... xs) noexcept { data = Native_vec_{xs...}; }
         __host__ __device__
diff --git a/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp b/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp
index 03e8158a4e..74bf932b3a 100644
--- a/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp
+++ b/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp
@@ -40,6 +40,7 @@ using namespace std;
 template<
     typename V,
     Enable_if_t<!is_integral<decltype(declval<V>().x)>{}>* = nullptr>
+__device__
 constexpr
 bool integer_unary_tests(const V&, const V&) {
     return true;
@@ -73,6 +74,7 @@ bool integer_unary_tests(V& f1, V& f2) {
 template<
     typename V,
     Enable_if_t<!is_integral<decltype(declval<V>().x)>{}>* = nullptr>
+__device__
 constexpr
 bool integer_binary_tests(const V&, const V&, const V&) {
     return true;