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
b9bf7fb8
"vscode:/vscode.git/clone" did not exist on "fad4eefa7ed52bae50617eae9a60e3719f72d604"
Commit
b9bf7fb8
authored
May 18, 2023
by
Rostyslav Geyyer
Browse files
Add some constexpr
parent
a30a0128
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
14 additions
and
10 deletions
+14
-10
include/ck/utility/f8_utils.hpp
include/ck/utility/f8_utils.hpp
+14
-10
No files found.
include/ck/utility/f8_utils.hpp
View file @
b9bf7fb8
...
@@ -42,25 +42,28 @@ __host__ __device__ f8_t cast_to_f8(float x, uint32_t rng)
...
@@ -42,25 +42,28 @@ __host__ __device__ f8_t cast_to_f8(float x, uint32_t rng)
sign
=
head
>>
(
we_f32
+
wm_f32
);
sign
=
head
>>
(
we_f32
+
wm_f32
);
uint32_t
signed_inf
=
(
sign
<<
(
we_f8
+
wm_f8
))
+
(((
1
<<
we_f8
)
-
1
)
<<
wm_f8
);
uint32_t
signed_inf
=
(
sign
<<
(
we_f8
+
wm_f8
))
+
(((
1
<<
we_f8
)
-
1
)
<<
wm_f8
);
uint32_t
drop_mask
=
(
1
<<
(
wm_f32
-
wm_f8
))
-
1
;
int
max_exp
;
int
exp_low_cutoff
;
if
(
negative_zero_nan
)
if
constexpr
(
negative_zero_nan
)
{
{
if
((
x_bitwise
&
0x7F800000
)
==
0x7F800000
)
if
((
x_bitwise
&
0x7F800000
)
==
0x7F800000
)
return
0x80
;
return
0x80
;
max_exp
=
(
1
<<
we_f8
)
-
1
;
exp_low_cutoff
=
0x80
-
(
1
<<
(
we_f8
-
1
));
}
}
else
else
{
{
if
((
x_bitwise
&
0x7F800000
)
==
0x7F800000
)
if
((
x_bitwise
&
0x7F800000
)
==
0x7F800000
)
return
signed_inf
+
(
mantissa
!=
0
?
1
:
0
);
return
signed_inf
+
(
mantissa
!=
0
?
1
:
0
);
max_exp
=
(
1
<<
we_f8
)
-
2
;
exp_low_cutoff
=
0x80
-
(
1
<<
(
we_f8
-
1
))
+
1
;
}
}
if
(
x_bitwise
==
0
)
if
(
x_bitwise
==
0
)
return
0
;
return
0
;
uint32_t
drop_mask
=
(
1
<<
(
wm_f32
-
wm_f8
))
-
1
;
const
int
max_exp
=
(
1
<<
we_f8
)
-
(
negative_zero_nan
?
1
:
2
);
const
int
exp_low_cutoff
=
0x80
-
(
1
<<
(
we_f8
-
1
))
+
1
-
(
negative_zero_nan
?
1
:
0
);
exponent
-=
exp_low_cutoff
-
1
;
exponent
-=
exp_low_cutoff
-
1
;
if
(
exponent
<=
0
)
if
(
exponent
<=
0
)
drop_mask
=
(
1
<<
(
wm_f32
-
wm_f8
+
1
-
exponent
))
-
1
;
drop_mask
=
(
1
<<
(
wm_f32
-
wm_f8
+
1
-
exponent
))
-
1
;
...
@@ -134,10 +137,14 @@ __host__ __device__ float cast_from_f8(f8_t x)
...
@@ -134,10 +137,14 @@ __host__ __device__ float cast_from_f8(f8_t x)
uint32_t
mantissa
=
x
&
((
1
<<
wm_f8
)
-
1
);
uint32_t
mantissa
=
x
&
((
1
<<
wm_f8
)
-
1
);
int
exponent
=
(
x
&
0x7F
)
>>
wm_f8
;
int
exponent
=
(
x
&
0x7F
)
>>
wm_f8
;
if
(
negative_zero_nan
)
int
exp_low_cutoff
;
uint32_t
retval
;
if
constexpr
(
negative_zero_nan
)
{
{
if
(
x
==
0x80
)
if
(
x
==
0x80
)
return
fNaN
;
return
fNaN
;
exp_low_cutoff
=
(
1
<<
(
we_f32
-
1
))
-
(
1
<<
(
we_f8
-
1
));
}
}
else
else
{
{
...
@@ -145,12 +152,9 @@ __host__ __device__ float cast_from_f8(f8_t x)
...
@@ -145,12 +152,9 @@ __host__ __device__ float cast_from_f8(f8_t x)
return
fNeg0
;
return
fNeg0
;
if
(
exponent
==
((
1
<<
we_f8
)
-
1
))
if
(
exponent
==
((
1
<<
we_f8
)
-
1
))
return
(
mantissa
==
0
)
?
(
sign
?
fNegInf
:
fInf
)
:
fNaN
;
return
(
mantissa
==
0
)
?
(
sign
?
fNegInf
:
fInf
)
:
fNaN
;
exp_low_cutoff
=
(
1
<<
(
we_f32
-
1
))
-
(
1
<<
(
we_f8
-
1
))
+
1
;
}
}
uint32_t
retval
;
const
int
exp_low_cutoff
=
(
1
<<
(
we_f32
-
1
))
-
(
1
<<
(
we_f8
-
1
))
+
1
-
(
negative_zero_nan
?
1
:
0
);
// subnormal input
// subnormal input
if
(
exponent
==
0
)
if
(
exponent
==
0
)
{
{
...
...
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