From 6005a9d8d60a3fb6c772150235f738fc9c7b614e 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: 3e2755adabdce7fc30b0a6af85c4b3b1fc1ed4ca]
---
 .../include/hip/hcc_detail/hip_vector_types.h | 77 +++++++++++++------
 .../src/deviceLib/hipVectorTypesDevice.cpp    |  2 +
 2 files changed, 55 insertions(+), 24 deletions(-)

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..60d53f4885 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__
@@ -682,32 +683,60 @@ __MAKE_VECTOR_TYPE__(longlong, long long);
 __MAKE_VECTOR_TYPE__(float, float);
 __MAKE_VECTOR_TYPE__(double, double);
 
-#define DECLOP_MAKE_ONE_COMPONENT(comp, type) \
-    __device__ __host__ \
-    static \
-    inline \
-    type make_##type(comp x) { type r = {x}; return r; }
+#if defined(__cplusplus)
+    #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \
+        __device__ __host__ \
+        static \
+        inline \
+        type make_##type(comp x) { return type{x}; }
 
-#define DECLOP_MAKE_TWO_COMPONENT(comp, type) \
-    __device__ __host__ \
-    static \
-    inline \
-    type make_##type(comp x, comp y) { type r = {x, y}; return r; }
+    #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \
+        __device__ __host__ \
+        static \
+        inline \
+        type make_##type(comp x, comp y) { return type{x, y}; }
 
-#define DECLOP_MAKE_THREE_COMPONENT(comp, type) \
-    __device__ __host__ \
-    static \
-    inline \
-    type make_##type(comp x, comp y, comp z) { type r = {x, y, z}; return r; }
+    #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \
+        __device__ __host__ \
+        static \
+        inline \
+        type make_##type(comp x, comp y, comp z) { return type{x, y, z}; }
 
-#define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \
-    __device__ __host__ \
-    static \
-    inline \
-    type make_##type(comp x, comp y, comp z, comp w) { \
-        type r = {x, y, z, w}; \
-        return r; \
-    }
+    #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \
+        __device__ __host__ \
+        static \
+        inline \
+        type make_##type(comp x, comp y, comp z, comp w) { \
+            return type{x, y, z, w}; \
+        }
+#else
+    #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \
+        __device__ __host__ \
+        static \
+        inline \
+        type make_##type(comp x) { type r = {x}; return r; }
+
+    #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \
+        __device__ __host__ \
+        static \
+        inline \
+        type make_##type(comp x, comp y) { type r = {x, y}; return r; }
+
+    #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \
+        __device__ __host__ \
+        static \
+        inline \
+        type make_##type(comp x, comp y, comp z) { type r = {x, y, z}; return r; }
+
+    #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \
+        __device__ __host__ \
+        static \
+        inline \
+        type make_##type(comp x, comp y, comp z, comp w) { \
+            type r = {x, y, z, w}; \
+            return r; \
+        }
+#endif
 
 DECLOP_MAKE_ONE_COMPONENT(unsigned char, uchar1);
 DECLOP_MAKE_TWO_COMPONENT(unsigned char, uchar2);
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;