From 2c568a554188b0ce1f944d0fd144e583018b33bd 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 [ROCm/clr commit: 1734eaead8183d34a11f3d4402362a37075850dc] --- .../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/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h index cf7058af2b..60d53f4885 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/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::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/clr/hipamd/tests/src/deviceLib/hipVectorTypesDevice.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipVectorTypesDevice.cpp index 03e8158a4e..74bf932b3a 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipVectorTypesDevice.cpp +++ b/projects/clr/hipamd/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 79fb5e969446a71a1bce0d1793d6d8019406a4d9 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 [ROCm/clr commit: 31efcc17d7e3584af43d075664d1ac6ebdc447b3] --- projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h | 3 ++- .../clr/hipamd/tests/src/deviceLib/hipVectorTypesDevice.cpp | 2 ++ 2 files changed, 4 insertions(+), 1 deletion(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h index cf7058af2b..371dfad189 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/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::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/clr/hipamd/tests/src/deviceLib/hipVectorTypesDevice.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipVectorTypesDevice.cpp index 03e8158a4e..74bf932b3a 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipVectorTypesDevice.cpp +++ b/projects/clr/hipamd/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 766211b400f2a327b1ad1fd50ab78148590abae3 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. [ROCm/clr commit: 14281c0a404436844960f84e0fa5ddad10c3837b] --- .../include/hip/hcc_detail/hip_vector_types.h | 73 ++++++------------- 1 file changed, 23 insertions(+), 50 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h index 60d53f4885..dc93b5697c 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/projects/clr/hipamd/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 c9f09d9e5d732c257914567be20d83ad01b30373 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. [ROCm/clr commit: 1ce56078f94b5203d19ae78bf24ad7a2de215c63] --- .../include/hip/hcc_detail/hip_vector_types.h | 74 ++++++------------- 1 file changed, 23 insertions(+), 51 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h index 60d53f4885..371dfad189 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/projects/clr/hipamd/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 8b6a48d0d1770366cc777a92a02a2253b4b4e841 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 [ROCm/clr commit: 252776d4aa969e40d08c79b66ce834e1c1b63aa4] --- projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h index dc93b5697c..7b05c7db3d 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/projects/clr/hipamd/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 48de002fb4e2a3fa400d648abc6b34baef9ff8e5 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. [ROCm/clr commit: 7020d118dace87554597aac7f17bcff4d22bbd09] --- .../hipamd/tests/src/deviceLib/hipVectorTypes.cpp | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipVectorTypes.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipVectorTypes.cpp index 734878b516..44681e3b05 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipVectorTypes.cpp +++ b/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().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;