From 3e2755adabdce7fc30b0a6af85c4b3b1fc1ed4ca Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 17 Jul 2018 17:17:48 +0100 Subject: [PATCH 1/6] Constrain variadic constructor to ranks > 1 and add missing `__device__`s --- include/hip/hcc_detail/hip_vector_types.h | 77 ++++++++++++++------ tests/src/deviceLib/hipVectorTypesDevice.cpp | 2 + 2 files changed, 55 insertions(+), 24 deletions(-) diff --git a/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h index cf7058af2b..60d53f4885 100644 --- a/include/hip/hcc_detail/hip_vector_types.h +++ b/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::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/tests/src/deviceLib/hipVectorTypesDevice.cpp b/tests/src/deviceLib/hipVectorTypesDevice.cpp index 03e8158a4e..74bf932b3a 100644 --- a/tests/src/deviceLib/hipVectorTypesDevice.cpp +++ b/tests/src/deviceLib/hipVectorTypesDevice.cpp @@ -40,6 +40,7 @@ using namespace std; template< typename V, Enable_if_t().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().x)>{}>* = nullptr> +__device__ constexpr bool integer_binary_tests(const V&, const V&, const V&) { return true; From 96aa6e22ee2197ecdad16c882dabe1e0298e1dc4 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 17 Jul 2018 17:17:48 +0100 Subject: [PATCH 2/6] Constrain variadic constructor to ranks > 1 and add missing `__device__`s --- include/hip/hcc_detail/hip_vector_types.h | 3 ++- tests/src/deviceLib/hipVectorTypesDevice.cpp | 2 ++ 2 files changed, 4 insertions(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h index cf7058af2b..371dfad189 100644 --- a/include/hip/hcc_detail/hip_vector_types.h +++ b/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::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/tests/src/deviceLib/hipVectorTypesDevice.cpp b/tests/src/deviceLib/hipVectorTypesDevice.cpp index 03e8158a4e..74bf932b3a 100644 --- a/tests/src/deviceLib/hipVectorTypesDevice.cpp +++ b/tests/src/deviceLib/hipVectorTypesDevice.cpp @@ -40,6 +40,7 @@ using namespace std; template< typename V, Enable_if_t().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().x)>{}>* = nullptr> +__device__ constexpr bool integer_binary_tests(const V&, const V&, const V&) { return true; From 73d0091328a000a52249bb2d63959e918b6c7d9c Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 17 Jul 2018 17:21:10 +0100 Subject: [PATCH 3/6] Not yet. --- include/hip/hcc_detail/hip_vector_types.h | 73 +++++++---------------- 1 file changed, 23 insertions(+), 50 deletions(-) diff --git a/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h index 60d53f4885..dc93b5697c 100644 --- a/include/hip/hcc_detail/hip_vector_types.h +++ b/include/hip/hcc_detail/hip_vector_types.h @@ -683,60 +683,33 @@ __MAKE_VECTOR_TYPE__(longlong, long long); __MAKE_VECTOR_TYPE__(float, float); __MAKE_VECTOR_TYPE__(double, double); -#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) { return type{x, y}; } +#define DECLOP_MAKE_ONE_COMPONENT(comp, type) \ + __device__ __host__ \ + static \ + inline \ + type make_##type(comp x) { type r = {x}; 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_TWO_COMPONENT(comp, type) \ + __device__ __host__ \ + static \ + inline \ + type make_##type(comp x, comp y) { type r = {x, y}; 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_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_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 +#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; \ + } DECLOP_MAKE_ONE_COMPONENT(unsigned char, uchar1); DECLOP_MAKE_TWO_COMPONENT(unsigned char, uchar2); From fe9859b79f7dcb5b57daa95cbf6dc410adcf1398 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 17 Jul 2018 17:24:48 +0100 Subject: [PATCH 4/6] Not yet. --- include/hip/hcc_detail/hip_vector_types.h | 74 +++++++---------------- 1 file changed, 23 insertions(+), 51 deletions(-) diff --git a/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h index 60d53f4885..371dfad189 100644 --- a/include/hip/hcc_detail/hip_vector_types.h +++ b/include/hip/hcc_detail/hip_vector_types.h @@ -683,60 +683,32 @@ __MAKE_VECTOR_TYPE__(longlong, long long); __MAKE_VECTOR_TYPE__(float, float); __MAKE_VECTOR_TYPE__(double, double); -#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_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) { return type{x, y}; } +#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) { return type{x, y, z}; } +#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) { \ - 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 +#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; \ + } DECLOP_MAKE_ONE_COMPONENT(unsigned char, uchar1); DECLOP_MAKE_TWO_COMPONENT(unsigned char, uchar2); From 487bb6f2efb8ca6c908d2e33fe43bba45868f48b Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 17 Jul 2018 17:28:48 +0100 Subject: [PATCH 5/6] Update hip_vector_types.h --- include/hip/hcc_detail/hip_vector_types.h | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h index dc93b5697c..7b05c7db3d 100644 --- a/include/hip/hcc_detail/hip_vector_types.h +++ b/include/hip/hcc_detail/hip_vector_types.h @@ -683,7 +683,6 @@ __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 \ @@ -771,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 \ No newline at end of file +#endif From 05ecc3582583b81e0dda1db58a720d478e63cbfd Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 17 Jul 2018 19:33:13 +0100 Subject: [PATCH 6/6] And a few more oversights. --- tests/src/deviceLib/hipVectorTypes.cpp | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/tests/src/deviceLib/hipVectorTypes.cpp b/tests/src/deviceLib/hipVectorTypes.cpp index 734878b516..44681e3b05 100644 --- a/tests/src/deviceLib/hipVectorTypes.cpp +++ b/tests/src/deviceLib/hipVectorTypes.cpp @@ -38,18 +38,23 @@ THE SOFTWARE. using namespace std; -bool integer_unary_tests(...) { +template< + typename V, + Enable_if_t().x)>{}>* = nullptr> +bool integer_unary_tests(V&, V&) { return true; } -bool integer_binary_tests(...) { +template< + typename V, + Enable_if_t().x)>{}>* = nullptr> +bool integer_binary_tests(V&, V&, V&...) { return true; } template< typename V, Enable_if_t().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().x)>{}>* = nullptr> -__device__ bool integer_binary_tests(V& f1, V& f2, V& f3) { f3 = f1 % f2; if (f3 != V{0}) return false;