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
Commits
f07a74d1
Commit
f07a74d1
authored
May 08, 2023
by
Rostyslav Geyyer
Browse files
Format
parent
21481b44
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
6 additions
and
7 deletions
+6
-7
include/ck/utility/data_type.hpp
include/ck/utility/data_type.hpp
+6
-7
No files found.
include/ck/utility/data_type.hpp
View file @
f07a74d1
...
@@ -1053,7 +1053,7 @@ template <>
...
@@ -1053,7 +1053,7 @@ template <>
inline
__host__
__device__
f8_t
type_convert
<
f8_t
,
float
>
(
float
x
)
inline
__host__
__device__
f8_t
type_convert
<
f8_t
,
float
>
(
float
x
)
{
{
constexpr
bool
negative_zero_nan
=
true
;
constexpr
bool
negative_zero_nan
=
true
;
constexpr
bool
clip
=
true
;
constexpr
bool
clip
=
true
;
// fp8 exponent/mantissa layout
// fp8 exponent/mantissa layout
constexpr
int
we
=
4
;
constexpr
int
we
=
4
;
...
@@ -1088,14 +1088,13 @@ inline __host__ __device__ f8_t type_convert<f8_t, float>(float x)
...
@@ -1088,14 +1088,13 @@ inline __host__ __device__ f8_t type_convert<f8_t, float>(float x)
if
((
_x
&
0x7F800000
)
==
0x7F800000
)
if
((
_x
&
0x7F800000
)
==
0x7F800000
)
return
signed_inf
+
(
mantissa
!=
0
?
1
:
0
);
return
signed_inf
+
(
mantissa
!=
0
?
1
:
0
);
}
}
if
(
_x
==
0
)
if
(
_x
==
0
)
return
0
;
return
0
;
uint32_t
drop_mask
=
(
1
<<
(
mfmt
-
wm
))
-
1
;
uint32_t
drop_mask
=
(
1
<<
(
mfmt
-
wm
))
-
1
;
const
int
max_exp
=
(
1
<<
we
)
-
(
negative_zero_nan
?
1
:
2
);
const
int
max_exp
=
(
1
<<
we
)
-
(
negative_zero_nan
?
1
:
2
);
const
int
exp_low_cutoff
=
const
int
exp_low_cutoff
=
128
-
(
1
<<
(
we
-
1
))
+
1
-
(
negative_zero_nan
?
1
:
0
);
128
-
(
1
<<
(
we
-
1
))
+
1
-
(
negative_zero_nan
?
1
:
0
);
exponent
-=
exp_low_cutoff
-
1
;
exponent
-=
exp_low_cutoff
-
1
;
if
(
exponent
<=
0
)
if
(
exponent
<=
0
)
...
@@ -1169,7 +1168,7 @@ inline __host__ __device__ float type_convert<float, f8_t>(f8_t x)
...
@@ -1169,7 +1168,7 @@ inline __host__ __device__ float type_convert<float, f8_t>(f8_t x)
uint32_t
sign
=
x
>>
7
;
uint32_t
sign
=
x
>>
7
;
uint32_t
mantissa
=
x
&
((
1
<<
wm
)
-
1
);
uint32_t
mantissa
=
x
&
((
1
<<
wm
)
-
1
);
int
exponent
=
(
x
&
0x7F
)
>>
wm
;
int
exponent
=
(
x
&
0x7F
)
>>
wm
;
if
(
negative_zero_nan
)
if
(
negative_zero_nan
)
{
{
if
(
x
==
0x80
)
if
(
x
==
0x80
)
...
...
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