Merge pull request #576 from ROCm-Developer-Tools/fix_missing_vector_bits

Constrain variadic constructor to ranks > 1 and add missing `__device__`s

[ROCm/clr commit: 6dc286b564]

projects/clr/hipamd/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__
@@ -769,4 +770,4 @@ DECLOP_MAKE_TWO_COMPONENT(signed long, longlong2);
 DECLOP_MAKE_THREE_COMPONENT(signed long, longlong3);
 DECLOP_MAKE_FOUR_COMPONENT(signed long, longlong4);
 
-#endif
+#endif

projects/clr/hipamd/tests/src/deviceLib/hipVectorTypes.cpp

@@ -38,18 +38,23 @@ THE SOFTWARE.
 
 using namespace std;
 
-bool integer_unary_tests(...) {
+template<
+    typename V,
+    Enable_if_t<!is_integral<decltype(declval<V>().x)>{}>* = nullptr>
+bool integer_unary_tests(V&, V&) {
     return true;
 }
 
-bool integer_binary_tests(...) {
+template<
+    typename V,
+    Enable_if_t<!is_integral<decltype(declval<V>().x)>{}>* = nullptr>
+bool integer_binary_tests(V&, V&, V&...) {
     return true;
 }
 
 template<
     typename V,
     Enable_if_t<is_integral<decltype(declval<V>().x)>{}>* = nullptr>
-__device__
 bool integer_unary_tests(V& f1, V& f2) {
     f1 %= f2;
     if (f1 != V{0}) return false;
@@ -71,7 +76,6 @@ bool integer_unary_tests(V& f1, V& f2) {
 template<
     typename V,
     Enable_if_t<is_integral<decltype(declval<V>().x)>{}>* = nullptr>
-__device__
 bool integer_binary_tests(V& f1, V& f2, V& f3) {
     f3 = f1 % f2;
     if (f3 != V{0}) return false;

projects/clr/hipamd/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;