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
fd2e6309
"vscode:/vscode.git/clone" did not exist on "0071478d9e770a81f1c79dca38bf788b5dd390fe"
Commit
fd2e6309
authored
May 15, 2023
by
Rostyslav Geyyer
Browse files
Eliminate magic numbers
parent
f2cf634f
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
5 additions
and
5 deletions
+5
-5
include/ck/utility/f8_utils.hpp
include/ck/utility/f8_utils.hpp
+5
-5
No files found.
include/ck/utility/f8_utils.hpp
View file @
fd2e6309
...
@@ -25,7 +25,7 @@ __host__ __device__ f8_t cast_to_f8(float x, uint32_t rng)
...
@@ -25,7 +25,7 @@ __host__ __device__ f8_t cast_to_f8(float x, uint32_t rng)
constexpr
int
wm_f8
=
3
;
constexpr
int
wm_f8
=
3
;
// fp32 exponent/mantissa layout
// fp32 exponent/mantissa layout
//
constexpr int we_f32 = 8;
constexpr
int
we_f32
=
8
;
constexpr
int
wm_f32
=
23
;
constexpr
int
wm_f32
=
23
;
uint32_t
x_bitwise
;
uint32_t
x_bitwise
;
...
@@ -38,8 +38,8 @@ __host__ __device__ f8_t cast_to_f8(float x, uint32_t rng)
...
@@ -38,8 +38,8 @@ __host__ __device__ f8_t cast_to_f8(float x, uint32_t rng)
head
=
x_bitwise
&
0xFF800000
;
head
=
x_bitwise
&
0xFF800000
;
mantissa
=
x_bitwise
&
0x7FFFFF
;
mantissa
=
x_bitwise
&
0x7FFFFF
;
exponent
=
(
head
>>
2
3
)
&
0xFF
;
exponent
=
(
head
>>
wm_f3
2
)
&
0xFF
;
sign
=
head
>>
31
;
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
);
...
@@ -99,9 +99,9 @@ __host__ __device__ f8_t cast_to_f8(float x, uint32_t rng)
...
@@ -99,9 +99,9 @@ __host__ __device__ f8_t cast_to_f8(float x, uint32_t rng)
}
}
}
}
if
(
exponent
==
0
&&
mantissa
==
0
)
if
(
exponent
==
0
&&
mantissa
==
0
)
return
negative_zero_nan
?
0
:
(
sign
<<
7
);
return
negative_zero_nan
?
0
:
(
sign
<<
(
we_f8
+
wm_f8
)
);
mantissa
&=
(
1
<<
wm_f8
)
-
1
;
mantissa
&=
(
1
<<
wm_f8
)
-
1
;
return
(
sign
<<
7
)
|
(
exponent
<<
wm_f8
)
|
mantissa
;
return
(
sign
<<
(
we_f8
+
wm_f8
)
)
|
(
exponent
<<
wm_f8
)
|
mantissa
;
}
}
// cast fp8 to fp32
// cast fp8 to fp32
...
...
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