Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in / Register
Toggle navigation
Menu
Open sidebar
gaoqiong
composable_kernel_ROCM
Commits
c0c1a7c1
Commit
c0c1a7c1
authored
Oct 08, 2024
by
Astha Rai
Browse files
added replicated functionality for standard header methods in utility files
parent
72c397c6
Changes
4
Hide whitespace changes
Inline
Side-by-side
Showing
4 changed files
with
130 additions
and
29 deletions
+130
-29
include/ck/utility/is_detected.hpp
include/ck/utility/is_detected.hpp
+8
-6
include/ck/utility/statically_indexed_array_multi_index.hpp
include/ck/utility/statically_indexed_array_multi_index.hpp
+17
-22
include/ck/utility/tuple_helper.hpp
include/ck/utility/tuple_helper.hpp
+1
-1
include/ck/utility/type.hpp
include/ck/utility/type.hpp
+104
-0
No files found.
include/ck/utility/is_detected.hpp
View file @
c0c1a7c1
...
@@ -3,20 +3,22 @@
...
@@ -3,20 +3,22 @@
#pragma once
#pragma once
#include "ck/utility/integral_constant.hpp"
namespace
ck
{
namespace
ck
{
namespace
detail
{
namespace
detail
{
template
<
class
Default
,
class
AlwaysVoid
,
template
<
class
...
>
class
Op
,
class
...
Args
>
template
<
class
Default
,
class
AlwaysVoid
,
template
<
class
...
>
class
Op
,
class
...
Args
>
struct
detector
struct
detector
{
{
using
value_t
=
std
::
false_type
;
using
value_t
=
integral_constant
<
bool
,
false
>
;
using
type
=
Default
;
using
type
=
Default
;
};
};
template
<
class
Default
,
template
<
class
...
>
class
Op
,
class
...
Args
>
template
<
class
Default
,
template
<
class
...
>
class
Op
,
class
...
Args
>
struct
detector
<
Default
,
std
::
void_t
<
Op
<
Args
...
>>
,
Op
,
Args
...
>
struct
detector
<
Default
,
ck
::
void_t
<
Op
<
Args
...
>>
,
Op
,
Args
...
>
{
{
using
value_t
=
std
::
true_type
;
using
value_t
=
integral_constant
<
bool
,
true
>
;
using
type
=
Op
<
Args
...
>
;
using
type
=
Op
<
Args
...
>
;
};
};
}
// namespace detail
}
// namespace detail
...
@@ -32,12 +34,12 @@ template <template <class...> class Op, class... Args>
...
@@ -32,12 +34,12 @@ template <template <class...> class Op, class... Args>
using
is_detected
=
typename
detail
::
detector
<
nonesuch
,
void
,
Op
,
Args
...
>::
value_t
;
using
is_detected
=
typename
detail
::
detector
<
nonesuch
,
void
,
Op
,
Args
...
>::
value_t
;
template
<
typename
T
>
template
<
typename
T
>
using
is_pack2_invocable_t
=
decltype
(
std
::
declval
<
T
&>
().
is_pack2_invocable
);
using
is_pack2_invocable_t
=
decltype
(
ck
::
declval
<
T
&>
().
is_pack2_invocable
);
template
<
typename
T
>
template
<
typename
T
>
using
is_pack4_invocable_t
=
decltype
(
std
::
declval
<
T
&>
().
is_pack4_invocable
);
using
is_pack4_invocable_t
=
decltype
(
ck
::
declval
<
T
&>
().
is_pack4_invocable
);
template
<
typename
T
>
template
<
typename
T
>
using
is_pack8_invocable_t
=
decltype
(
std
::
declval
<
T
&>
().
is_pack8_invocable
);
using
is_pack8_invocable_t
=
decltype
(
ck
::
declval
<
T
&>
().
is_pack8_invocable
);
}
// namespace ck
}
// namespace ck
include/ck/utility/statically_indexed_array_multi_index.hpp
View file @
c0c1a7c1
...
@@ -35,10 +35,9 @@ __host__ __device__ constexpr auto to_multi_index(const T& x)
...
@@ -35,10 +35,9 @@ __host__ __device__ constexpr auto to_multi_index(const T& x)
// is the alias of the latter. This is because compiler cannot infer the NSize if
// is the alias of the latter. This is because compiler cannot infer the NSize if
// using MultiIndex<NSize>
// using MultiIndex<NSize>
// TODO: how to fix this?
// TODO: how to fix this?
template
<
template
<
typename
...
Ys
,
typename
...
Ys
,
typename
X
,
typename
X
,
enable_if_t
<!
ck
::
is_integral
<
X
>
::
value
&&
!
ck
::
is_floating_point
<
X
>::
value
,
bool
>
=
false
>
enable_if_t
<!
std
::
is_integral
<
X
>
::
value
&&
!
std
::
is_floating_point
<
X
>::
value
,
bool
>
=
false
>
__host__
__device__
constexpr
auto
operator
+=
(
Tuple
<
Ys
...
>&
y
,
const
X
&
x
)
__host__
__device__
constexpr
auto
operator
+=
(
Tuple
<
Ys
...
>&
y
,
const
X
&
x
)
{
{
static_assert
(
X
::
Size
()
==
sizeof
...(
Ys
),
"wrong! size not the same"
);
static_assert
(
X
::
Size
()
==
sizeof
...(
Ys
),
"wrong! size not the same"
);
...
@@ -47,10 +46,9 @@ __host__ __device__ constexpr auto operator+=(Tuple<Ys...>& y, const X& x)
...
@@ -47,10 +46,9 @@ __host__ __device__ constexpr auto operator+=(Tuple<Ys...>& y, const X& x)
return
y
;
return
y
;
}
}
template
<
template
<
typename
...
Ys
,
typename
...
Ys
,
typename
X
,
typename
X
,
enable_if_t
<!
ck
::
is_integral
<
X
>
::
value
&&
!
ck
::
is_floating_point
<
X
>::
value
,
bool
>
=
false
>
enable_if_t
<!
std
::
is_integral
<
X
>
::
value
&&
!
std
::
is_floating_point
<
X
>::
value
,
bool
>
=
false
>
__host__
__device__
constexpr
auto
operator
-=
(
Tuple
<
Ys
...
>&
y
,
const
X
&
x
)
__host__
__device__
constexpr
auto
operator
-=
(
Tuple
<
Ys
...
>&
y
,
const
X
&
x
)
{
{
static_assert
(
X
::
Size
()
==
sizeof
...(
Ys
),
"wrong! size not the same"
);
static_assert
(
X
::
Size
()
==
sizeof
...(
Ys
),
"wrong! size not the same"
);
...
@@ -59,10 +57,9 @@ __host__ __device__ constexpr auto operator-=(Tuple<Ys...>& y, const X& x)
...
@@ -59,10 +57,9 @@ __host__ __device__ constexpr auto operator-=(Tuple<Ys...>& y, const X& x)
return
y
;
return
y
;
}
}
template
<
template
<
typename
...
Xs
,
typename
...
Xs
,
typename
Y
,
typename
Y
,
enable_if_t
<!
ck
::
is_integral
<
Y
>
::
value
&&
!
ck
::
is_floating_point
<
Y
>::
value
,
bool
>
=
false
>
enable_if_t
<!
std
::
is_integral
<
Y
>
::
value
&&
!
std
::
is_floating_point
<
Y
>::
value
,
bool
>
=
false
>
__host__
__device__
constexpr
auto
operator
+
(
const
Tuple
<
Xs
...
>&
x
,
const
Y
&
y
)
__host__
__device__
constexpr
auto
operator
+
(
const
Tuple
<
Xs
...
>&
x
,
const
Y
&
y
)
{
{
static_assert
(
Y
::
Size
()
==
sizeof
...(
Xs
),
"wrong! size not the same"
);
static_assert
(
Y
::
Size
()
==
sizeof
...(
Xs
),
"wrong! size not the same"
);
...
@@ -73,10 +70,9 @@ __host__ __device__ constexpr auto operator+(const Tuple<Xs...>& x, const Y& y)
...
@@ -73,10 +70,9 @@ __host__ __device__ constexpr auto operator+(const Tuple<Xs...>& x, const Y& y)
return
r
;
return
r
;
}
}
template
<
template
<
typename
...
Xs
,
typename
...
Xs
,
typename
Y
,
typename
Y
,
enable_if_t
<!
ck
::
is_integral
<
Y
>
::
value
&&
!
ck
::
is_floating_point
<
Y
>::
value
,
bool
>
=
false
>
enable_if_t
<!
std
::
is_integral
<
Y
>
::
value
&&
!
std
::
is_floating_point
<
Y
>::
value
,
bool
>
=
false
>
__host__
__device__
constexpr
auto
operator
-
(
const
Tuple
<
Xs
...
>&
x
,
const
Y
&
y
)
__host__
__device__
constexpr
auto
operator
-
(
const
Tuple
<
Xs
...
>&
x
,
const
Y
&
y
)
{
{
static_assert
(
Y
::
Size
()
==
sizeof
...(
Xs
),
"wrong! size not the same"
);
static_assert
(
Y
::
Size
()
==
sizeof
...(
Xs
),
"wrong! size not the same"
);
...
@@ -87,10 +83,9 @@ __host__ __device__ constexpr auto operator-(const Tuple<Xs...>& x, const Y& y)
...
@@ -87,10 +83,9 @@ __host__ __device__ constexpr auto operator-(const Tuple<Xs...>& x, const Y& y)
return
r
;
return
r
;
}
}
template
<
template
<
typename
...
Xs
,
typename
...
Xs
,
typename
Y
,
typename
Y
,
enable_if_t
<!
ck
::
is_integral
<
Y
>
::
value
&&
!
ck
::
is_floating_point
<
Y
>::
value
,
bool
>
=
false
>
enable_if_t
<!
std
::
is_integral
<
Y
>
::
value
&&
!
std
::
is_floating_point
<
Y
>::
value
,
bool
>
=
false
>
__host__
__device__
constexpr
auto
operator
*
(
const
Tuple
<
Xs
...
>&
x
,
const
Y
&
y
)
__host__
__device__
constexpr
auto
operator
*
(
const
Tuple
<
Xs
...
>&
x
,
const
Y
&
y
)
{
{
static_assert
(
Y
::
Size
()
==
sizeof
...(
Xs
),
"wrong! size not the same"
);
static_assert
(
Y
::
Size
()
==
sizeof
...(
Xs
),
"wrong! size not the same"
);
...
@@ -104,7 +99,7 @@ __host__ __device__ constexpr auto operator*(const Tuple<Xs...>& x, const Y& y)
...
@@ -104,7 +99,7 @@ __host__ __device__ constexpr auto operator*(const Tuple<Xs...>& x, const Y& y)
// MultiIndex = scalar * MultiIndex
// MultiIndex = scalar * MultiIndex
template
<
typename
...
Xs
,
template
<
typename
...
Xs
,
typename
Y
,
typename
Y
,
enable_if_t
<
std
::
is_integral
<
Y
>
::
value
||
std
::
is_floating_point
<
Y
>::
value
,
bool
>
=
false
>
enable_if_t
<
ck
::
is_integral
<
Y
>
::
value
||
ck
::
is_floating_point
<
Y
>::
value
,
bool
>
=
false
>
__host__
__device__
constexpr
auto
operator
*
(
Y
a
,
const
Tuple
<
Xs
...
>&
x
)
__host__
__device__
constexpr
auto
operator
*
(
Y
a
,
const
Tuple
<
Xs
...
>&
x
)
{
{
constexpr
index_t
NSize
=
sizeof
...(
Xs
);
constexpr
index_t
NSize
=
sizeof
...(
Xs
);
...
@@ -117,7 +112,7 @@ __host__ __device__ constexpr auto operator*(Y a, const Tuple<Xs...>& x)
...
@@ -117,7 +112,7 @@ __host__ __device__ constexpr auto operator*(Y a, const Tuple<Xs...>& x)
// MultiIndex = MultiIndex * scalar
// MultiIndex = MultiIndex * scalar
template
<
typename
...
Xs
,
template
<
typename
...
Xs
,
typename
Y
,
typename
Y
,
enable_if_t
<
std
::
is_integral
<
Y
>
::
value
||
std
::
is_floating_point
<
Y
>::
value
,
bool
>
=
false
>
enable_if_t
<
ck
::
is_integral
<
Y
>
::
value
||
ck
::
is_floating_point
<
Y
>::
value
,
bool
>
=
false
>
__host__
__device__
constexpr
auto
operator
*
(
const
Tuple
<
Xs
...
>&
x
,
Y
a
)
__host__
__device__
constexpr
auto
operator
*
(
const
Tuple
<
Xs
...
>&
x
,
Y
a
)
{
{
return
a
*
x
;
return
a
*
x
;
...
...
include/ck/utility/tuple_helper.hpp
View file @
c0c1a7c1
...
@@ -161,7 +161,7 @@ __host__ __device__ constexpr auto TupleReduce(F&& f, const Tuple<Ts...>& tuple)
...
@@ -161,7 +161,7 @@ __host__ __device__ constexpr auto TupleReduce(F&& f, const Tuple<Ts...>& tuple)
#ifndef CK_CODE_GEN_RTC
#ifndef CK_CODE_GEN_RTC
template
<
typename
T
>
template
<
typename
T
>
using
is_tuple
=
decltype
(
std
::
declval
<
T
&>
().
IsTuple
());
using
is_tuple
=
decltype
(
ck
::
declval
<
T
&>
().
IsTuple
());
#endif
#endif
template
<
typename
...
Ts
>
template
<
typename
...
Ts
>
...
...
include/ck/utility/type.hpp
View file @
c0c1a7c1
...
@@ -119,6 +119,9 @@ constexpr T&& forward(typename remove_reference<T>::type&& t_) noexcept
...
@@ -119,6 +119,9 @@ constexpr T&& forward(typename remove_reference<T>::type&& t_) noexcept
template
<
typename
T
>
template
<
typename
T
>
T
&&
declval
()
noexcept
;
T
&&
declval
()
noexcept
;
template
<
typename
...
Ts
>
using
void_t
=
void
;
#else
#else
#include <utility>
#include <utility>
#include <type_traits>
#include <type_traits>
...
@@ -134,6 +137,7 @@ using std::remove_const;
...
@@ -134,6 +137,7 @@ using std::remove_const;
using
std
::
remove_cv
;
using
std
::
remove_cv
;
using
std
::
remove_pointer
;
using
std
::
remove_pointer
;
using
std
::
remove_reference
;
using
std
::
remove_reference
;
using
std
::
void_t
;
#endif
#endif
template
<
typename
X
,
typename
Y
>
template
<
typename
X
,
typename
Y
>
...
@@ -156,6 +160,106 @@ struct is_const<const X> : public integral_constant<bool, true>
...
@@ -156,6 +160,106 @@ struct is_const<const X> : public integral_constant<bool, true>
{
{
};
};
template
<
typename
X
>
struct
is_floating_point
:
public
integral_constant
<
bool
,
false
>
{
};
template
<
>
struct
is_floating_point
<
float
>
:
public
integral_constant
<
bool
,
true
>
{
};
template
<
>
struct
is_floating_point
<
double
>
:
public
integral_constant
<
bool
,
true
>
{
};
template
<
>
struct
is_floating_point
<
long
double
>
:
public
integral_constant
<
bool
,
true
>
{
};
template
<
typename
X
>
struct
is_integral
:
public
integral_constant
<
bool
,
false
>
{
};
template
<
>
struct
is_integral
<
int
>
:
public
integral_constant
<
bool
,
true
>
{
};
template
<
>
struct
is_integral
<
unsigned
int
>
:
public
integral_constant
<
bool
,
true
>
{
};
template
<
>
struct
is_integral
<
long
>
:
public
integral_constant
<
bool
,
true
>
{
};
template
<
>
struct
is_integral
<
unsigned
long
>
:
public
integral_constant
<
bool
,
true
>
{
};
template
<
>
struct
is_integral
<
short
>
:
public
integral_constant
<
bool
,
true
>
{
};
template
<
>
struct
is_integral
<
unsigned
short
>
:
public
integral_constant
<
bool
,
true
>
{
};
template
<
>
struct
is_integral
<
long
long
>
:
public
integral_constant
<
bool
,
true
>
{
};
template
<
>
struct
is_integral
<
unsigned
long
long
>
:
public
integral_constant
<
bool
,
true
>
{
};
template
<
>
struct
is_integral
<
char
>
:
public
integral_constant
<
bool
,
true
>
{
};
template
<
>
struct
is_integral
<
signed
char
>
:
public
integral_constant
<
bool
,
true
>
{
};
template
<
>
struct
is_integral
<
unsigned
char
>
:
public
integral_constant
<
bool
,
true
>
{
};
template
<
>
struct
is_integral
<
wchar_t
>
:
public
integral_constant
<
bool
,
true
>
{
};
template
<
>
struct
is_integral
<
char16_t
>
:
public
integral_constant
<
bool
,
true
>
{
};
template
<
>
struct
is_integral
<
char32_t
>
:
public
integral_constant
<
bool
,
true
>
{
};
template
<
>
struct
is_integral
<
bool
>
:
public
integral_constant
<
bool
,
true
>
{
};
template
<
typename
T
>
template
<
typename
T
>
inline
constexpr
bool
is_reference_v
=
is_reference
<
T
>::
value
;
inline
constexpr
bool
is_reference_v
=
is_reference
<
T
>::
value
;
...
...
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
.
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment