@@ -87,6 +87,9 @@ namespace internal
8787{
8888 template <camp::idx_t index, typename Type>
8989 struct tuple_storage {
90+ #if defined(__NVCC__)
91+ // NVCC up until at least 10.1 can't do aggregate initialization of a pack
92+ // of base classes, keep this around until that's fixed
9093 CAMP_HOST_DEVICE constexpr tuple_storage () : val(){};
9194
9295 /* Workaround for bug in hipcc compiler */
@@ -105,15 +108,8 @@ namespace internal
105108 : val(std::forward<T>(v))
106109 {
107110 }
111+ #endif
108112
109- CAMP_HOST_DEVICE constexpr const Type& get_inner () const noexcept
110- {
111- return val;
112- }
113-
114- CAMP_HOST_DEVICE CAMP_CONSTEXPR14 Type& get_inner () noexcept { return val; }
115-
116- public:
117113 Type val;
118114 };
119115
@@ -122,23 +118,34 @@ namespace internal
122118
123119 template <typename ... Types, camp::idx_t ... Indices>
124120 struct tuple_helper <camp::idx_seq<Indices...>, camp::list<Types...>>
125- : public internal::tuple_storage<Indices, Types>... {
121+ : public tuple_storage<Indices, Types>... {
122+
123+ #if !defined(__NVCC__)
124+ constexpr tuple_helper () = default;
125+ constexpr tuple_helper (tuple_helper const &) = default;
126+ constexpr tuple_helper (tuple_helper&&) = default;
127+
128+ #else
129+ // NOTE: this is to work around nvcc 9 series issues with incorrect
130+ // creation of defaulted constructors
126131 template <
127132 bool B = all_of<std::is_default_constructible<Types>::value...>::value,
128133 typename std::enable_if<B, void >::type* = nullptr >
129134 CAMP_HOST_DEVICE constexpr tuple_helper ()
135+ : tuple_storage<Indices, Types>()...
130136 {
131137 }
132138 CAMP_HOST_DEVICE constexpr tuple_helper (tuple_helper const & rhs)
133139 : tuple_storage<Indices, Types>(
134- rhs.tuple_storage<Indices, Types>::get_inner() )...
140+ rhs.tuple_storage<Indices, Types>::val )...
135141 {
136142 }
137143 CAMP_HOST_DEVICE constexpr tuple_helper (tuple_helper&& rhs)
138144 : tuple_storage<Indices, Types>(
139145 std::forward<Types>(rhs.tuple_storage<Indices, Types>::val))...
140146 {
141147 }
148+ #endif
142149
143150 /* Workaround for bug in hipcc compiler */
144151 // This likely causes issues when building with hip
@@ -152,7 +159,7 @@ namespace internal
152159
153160 template <typename ... Args>
154161 CAMP_HOST_DEVICE constexpr tuple_helper (Args&&... args)
155- : tuple_storage<Indices, Types>( std::forward<Args>(args)) ...
162+ : tuple_storage<Indices, Types>{ std::forward<Args>(args)} ...
156163 {
157164 }
158165
@@ -162,7 +169,7 @@ namespace internal
162169 template <typename RTuple>
163170 CAMP_HOST_DEVICE tuple_helper& operator =(const RTuple& rhs)
164171 {
165- return (camp::sink ((this ->tuple_storage <Indices, Types>::get_inner () =
172+ return (camp::sink ((this ->tuple_storage <Indices, Types>::val =
166173 get<Indices>(rhs))...),
167174 *this );
168175 }
@@ -220,12 +227,11 @@ struct tuple {
220227 -> tuple_ebt_t<T, Tuple>&;
221228
222229public:
223- // NOTE: __host__ __device__ on constructors causes warnings, and nothing else
224- // Constructors
225- template <
226- bool B = all_of<std::is_default_constructible<Elements>::value...>::value,
227- typename std::enable_if<B, void >::type* = nullptr >
228- CAMP_HOST_DEVICE constexpr tuple () : base()
230+ // NOTE: __host__ __device__ on constructors causes warnings, and nothing
231+ // else Constructors
232+ template <bool B = std::is_default_constructible<Base>::value,
233+ typename std::enable_if<B, void >::type* = nullptr >
234+ CAMP_HOST_DEVICE constexpr tuple () : base{}
229235 {
230236 }
231237
@@ -314,11 +320,10 @@ class tagged_tuple : public tuple<Elements...>
314320
315321
316322public:
317- // NOTE: __host__ __device__ on constructors causes warnings, and nothing else
318- // Constructors
319- template <
320- bool B = all_of<std::is_default_constructible<Elements>::value...>::value,
321- typename std::enable_if<B, void >::type* = nullptr >
323+ // NOTE: __host__ __device__ on constructors causes warnings, and nothing
324+ // else Constructors
325+ template <bool B = std::is_default_constructible<Base>::value,
326+ typename std::enable_if<B, void >::type* = nullptr >
322327 CAMP_HOST_DEVICE constexpr tagged_tuple () : base()
323328 {
324329 }
@@ -375,7 +380,7 @@ CAMP_HOST_DEVICE constexpr auto get(const Tuple& t) noexcept
375380{
376381 using internal::tpl_get_store;
377382 static_assert (tuple_size<Tuple>::value > index, " index out of range" );
378- return static_cast <tpl_get_store<Tuple, index> const &>(t.base ).get_inner () ;
383+ return static_cast <tpl_get_store<Tuple, index> const &>(t.base ).val ;
379384}
380385
381386template <camp::idx_t index, class Tuple >
@@ -384,7 +389,7 @@ CAMP_HOST_DEVICE constexpr auto get(Tuple& t) noexcept
384389{
385390 using internal::tpl_get_store;
386391 static_assert (tuple_size<Tuple>::value > index, " index out of range" );
387- return static_cast <tpl_get_store<Tuple, index>&>(t.base ).get_inner () ;
392+ return static_cast <tpl_get_store<Tuple, index>&>(t.base ).val ;
388393}
389394
390395// by type
@@ -397,8 +402,7 @@ CAMP_HOST_DEVICE constexpr auto get(const Tuple& t) noexcept
397402 static_assert (!std::is_same<camp::nil, index_type>::value,
398403 " invalid type index" );
399404
400- return static_cast <tpl_get_store<Tuple, index_type::value>&>(t.base )
401- .get_inner ();
405+ return static_cast <tpl_get_store<Tuple, index_type::value>&>(t.base ).val ;
402406}
403407
404408template <typename T, class Tuple >
@@ -409,8 +413,7 @@ CAMP_HOST_DEVICE constexpr auto get(Tuple& t) noexcept -> tuple_ebt_t<T, Tuple>&
409413 static_assert (!std::is_same<camp::nil, index_type>::value,
410414 " invalid type index" );
411415
412- return static_cast <tpl_get_store<Tuple, index_type::value>&>(t.base )
413- .get_inner ();
416+ return static_cast <tpl_get_store<Tuple, index_type::value>&>(t.base ).val ;
414417}
415418
416419template <typename ... Args>
0 commit comments