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
f714319e
Commit
f714319e
authored
Jan 13, 2025
by
Astha Rai
Browse files
removed and replicated standard header usage from data_type and type_convert files for fp8 changes
parent
3aec6f03
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
10 additions
and
10 deletions
+10
-10
include/ck/utility/data_type.hpp
include/ck/utility/data_type.hpp
+8
-8
include/ck/utility/type_convert.hpp
include/ck/utility/type_convert.hpp
+2
-2
No files found.
include/ck/utility/data_type.hpp
View file @
f714319e
...
@@ -327,7 +327,7 @@ struct vector_type<T, 2, typename ck::enable_if_t<is_native_type<T>()>>
...
@@ -327,7 +327,7 @@ struct vector_type<T, 2, typename ck::enable_if_t<is_native_type<T>()>>
};
};
template
<
typename
T
>
template
<
typename
T
>
struct
vector_type
<
T
,
3
,
typename
std
::
enable_if_t
<
is_native_type
<
T
>
()
>>
struct
vector_type
<
T
,
3
,
typename
ck
::
enable_if_t
<
is_native_type
<
T
>
()
>>
{
{
using
d1_t
=
T
;
using
d1_t
=
T
;
typedef
T
d2_t
__attribute__
((
ext_vector_type
(
2
)));
typedef
T
d2_t
__attribute__
((
ext_vector_type
(
2
)));
...
@@ -397,7 +397,7 @@ struct vector_type<T, 3, typename std::enable_if_t<is_native_type<T>()>>
...
@@ -397,7 +397,7 @@ struct vector_type<T, 3, typename std::enable_if_t<is_native_type<T>()>>
};
};
template
<
typename
T
>
template
<
typename
T
>
struct
vector_type
<
T
,
4
,
typename
std
::
enable_if_t
<
is_native_type
<
T
>
()
>>
struct
vector_type
<
T
,
4
,
typename
ck
::
enable_if_t
<
is_native_type
<
T
>
()
>>
{
{
using
d1_t
=
T
;
using
d1_t
=
T
;
typedef
T
d2_t
__attribute__
((
ext_vector_type
(
2
)));
typedef
T
d2_t
__attribute__
((
ext_vector_type
(
2
)));
...
@@ -467,7 +467,7 @@ struct vector_type<T, 4, typename std::enable_if_t<is_native_type<T>()>>
...
@@ -467,7 +467,7 @@ struct vector_type<T, 4, typename std::enable_if_t<is_native_type<T>()>>
};
};
template
<
typename
T
>
template
<
typename
T
>
struct
vector_type
<
T
,
5
,
typename
std
::
enable_if_t
<
is_native_type
<
T
>
()
>>
struct
vector_type
<
T
,
5
,
typename
ck
::
enable_if_t
<
is_native_type
<
T
>
()
>>
{
{
using
d1_t
=
T
;
using
d1_t
=
T
;
typedef
T
d4_t
__attribute__
((
ext_vector_type
(
4
)));
typedef
T
d4_t
__attribute__
((
ext_vector_type
(
4
)));
...
@@ -537,7 +537,7 @@ struct vector_type<T, 5, typename std::enable_if_t<is_native_type<T>()>>
...
@@ -537,7 +537,7 @@ struct vector_type<T, 5, typename std::enable_if_t<is_native_type<T>()>>
};
};
template
<
typename
T
>
template
<
typename
T
>
struct
vector_type
<
T
,
7
,
typename
std
::
enable_if_t
<
is_native_type
<
T
>
()
>>
struct
vector_type
<
T
,
7
,
typename
ck
::
enable_if_t
<
is_native_type
<
T
>
()
>>
{
{
using
d1_t
=
T
;
using
d1_t
=
T
;
typedef
T
d2_t
__attribute__
((
ext_vector_type
(
2
)));
typedef
T
d2_t
__attribute__
((
ext_vector_type
(
2
)));
...
@@ -619,7 +619,7 @@ struct vector_type<T, 7, typename std::enable_if_t<is_native_type<T>()>>
...
@@ -619,7 +619,7 @@ struct vector_type<T, 7, typename std::enable_if_t<is_native_type<T>()>>
};
};
template
<
typename
T
>
template
<
typename
T
>
struct
vector_type
<
T
,
8
,
typename
std
::
enable_if_t
<
is_native_type
<
T
>
()
>>
struct
vector_type
<
T
,
8
,
typename
ck
::
enable_if_t
<
is_native_type
<
T
>
()
>>
{
{
using
d1_t
=
T
;
using
d1_t
=
T
;
typedef
T
d2_t
__attribute__
((
ext_vector_type
(
2
)));
typedef
T
d2_t
__attribute__
((
ext_vector_type
(
2
)));
...
@@ -701,7 +701,7 @@ struct vector_type<T, 8, typename std::enable_if_t<is_native_type<T>()>>
...
@@ -701,7 +701,7 @@ struct vector_type<T, 8, typename std::enable_if_t<is_native_type<T>()>>
};
};
template
<
typename
T
>
template
<
typename
T
>
struct
vector_type
<
T
,
13
,
typename
std
::
enable_if_t
<
is_native_type
<
T
>
()
>>
struct
vector_type
<
T
,
13
,
typename
ck
::
enable_if_t
<
is_native_type
<
T
>
()
>>
{
{
using
d1_t
=
T
;
using
d1_t
=
T
;
typedef
T
d4_t
__attribute__
((
ext_vector_type
(
4
)));
typedef
T
d4_t
__attribute__
((
ext_vector_type
(
4
)));
...
@@ -783,7 +783,7 @@ struct vector_type<T, 13, typename std::enable_if_t<is_native_type<T>()>>
...
@@ -783,7 +783,7 @@ struct vector_type<T, 13, typename std::enable_if_t<is_native_type<T>()>>
};
};
template
<
typename
T
>
template
<
typename
T
>
struct
vector_type
<
T
,
16
,
typename
std
::
enable_if_t
<
is_native_type
<
T
>
()
>>
struct
vector_type
<
T
,
16
,
typename
ck
::
enable_if_t
<
is_native_type
<
T
>
()
>>
{
{
using
d1_t
=
T
;
using
d1_t
=
T
;
typedef
T
d2_t
__attribute__
((
ext_vector_type
(
2
)));
typedef
T
d2_t
__attribute__
((
ext_vector_type
(
2
)));
...
@@ -1388,7 +1388,7 @@ template <typename T, index_t N>
...
@@ -1388,7 +1388,7 @@ template <typename T, index_t N>
struct
non_native_vector_base
<
struct
non_native_vector_base
<
T
,
T
,
N
,
N
,
std
::
enable_if_t
<
sizeof
(
T
)
==
1
||
sizeof
(
T
)
==
2
||
sizeof
(
T
)
==
4
||
sizeof
(
T
)
==
8
>>
ck
::
enable_if_t
<
sizeof
(
T
)
==
1
||
sizeof
(
T
)
==
2
||
sizeof
(
T
)
==
4
||
sizeof
(
T
)
==
8
>>
{
{
using
data_t
=
typename
nnvb_data_t_selector
<
T
>::
type
;
// select data_t based on the size of T
using
data_t
=
typename
nnvb_data_t_selector
<
T
>::
type
;
// select data_t based on the size of T
static_assert
(
sizeof
(
T
)
==
sizeof
(
data_t
),
"non_native_vector_base storage size mismatch"
);
static_assert
(
sizeof
(
T
)
==
sizeof
(
data_t
),
"non_native_vector_base storage size mismatch"
);
...
...
include/ck/utility/type_convert.hpp
View file @
f714319e
...
@@ -219,7 +219,7 @@ inline __host__ __device__ f8_fnuz_t f8_convert_sr<f8_fnuz_t, half_t>(half_t x)
...
@@ -219,7 +219,7 @@ inline __host__ __device__ f8_fnuz_t f8_convert_sr<f8_fnuz_t, half_t>(half_t x)
constexpr
f8_rounding_mode
rm
=
f8_rounding_mode
::
stochastic
;
constexpr
f8_rounding_mode
rm
=
f8_rounding_mode
::
stochastic
;
constexpr
int
seed
=
1254739
;
constexpr
int
seed
=
1254739
;
uint32_t
rng
=
prand_generator
<
half_t
,
seed
>
(
reinterpret_cast
<
uintptr
_t
>
(
&
x
),
x
);
uint32_t
rng
=
prand_generator
<
half_t
,
seed
>
(
reinterpret_cast
<
size
_t
>
(
&
x
),
x
);
return
utils
::
cast_to_f8
<
half_t
,
return
utils
::
cast_to_f8
<
half_t
,
f8_fnuz_t
,
f8_fnuz_t
,
negative_zero_nan
,
negative_zero_nan
,
...
@@ -276,7 +276,7 @@ inline __host__ __device__ bf8_fnuz_t f8_convert_sr<bf8_fnuz_t, half_t>(half_t x
...
@@ -276,7 +276,7 @@ inline __host__ __device__ bf8_fnuz_t f8_convert_sr<bf8_fnuz_t, half_t>(half_t x
constexpr
f8_rounding_mode
rm
=
f8_rounding_mode
::
stochastic
;
constexpr
f8_rounding_mode
rm
=
f8_rounding_mode
::
stochastic
;
constexpr
int
seed
=
1254739
;
constexpr
int
seed
=
1254739
;
uint32_t
rng
=
prand_generator
<
half_t
,
seed
>
(
reinterpret_cast
<
uintptr
_t
>
(
&
x
),
x
);
uint32_t
rng
=
prand_generator
<
half_t
,
seed
>
(
reinterpret_cast
<
size
_t
>
(
&
x
),
x
);
return
utils
::
cast_to_f8
<
half_t
,
return
utils
::
cast_to_f8
<
half_t
,
bf8_fnuz_t
,
bf8_fnuz_t
,
negative_zero_nan
,
negative_zero_nan
,
...
...
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