Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Device kernels using FP8 #2510

Merged
merged 107 commits into from
Dec 6, 2023
Merged
Changes from 1 commit
Commits
Show all changes
107 commits
Select commit Hold shift + click to select a range
df7f8a3
changes for the FP8 ref implementation
umangyadav Nov 9, 2023
9bc1828
cppcheck fixes
umangyadav Nov 9, 2023
155a2b1
move FNUZ as template parameter
umangyadav Nov 10, 2023
d9f11e3
Fix numeric limits
umangyadav Nov 10, 2023
4e9d51f
Working FNUZ and FN
umangyadav Nov 10, 2023
7639c28
use float equal
umangyadav Nov 10, 2023
a6372c5
add test for fp8e5m2
umangyadav Nov 10, 2023
439ea40
add test for fp8e5m2fnuz
umangyadav Nov 10, 2023
183db78
refactor add some comments
umangyadav Nov 10, 2023
ab653af
Review updates
umangyadav Nov 13, 2023
8319e01
Fix tidy
umangyadav Nov 14, 2023
9ee0418
Fix test failure
umangyadav Nov 14, 2023
355e4f6
fix isfinite
umangyadav Nov 14, 2023
ba471f4
Merge remote-tracking branch 'origin/develop' into ref_fp8
umangyadav Nov 14, 2023
6aec703
fix test for neg inf
umangyadav Nov 14, 2023
12aac37
fix warning
umangyadav Nov 14, 2023
6009232
add tests
umangyadav Nov 14, 2023
03f7139
Fix tests
umangyadav Nov 14, 2023
1e220c0
add stringstream tests
umangyadav Nov 14, 2023
a83e9dc
Remove clang diagnostics
umangyadav Nov 15, 2023
dfb35a6
Merge remote-tracking branch 'origin/develop' into ref_fp8
umangyadav Nov 15, 2023
26956f1
Remove NOLINTS
umangyadav Nov 15, 2023
269ce6d
Bugfixes and additional tests
umangyadav Nov 16, 2023
6414ee3
Fix undoing
umangyadav Nov 16, 2023
cd26ada
Handle underflow case separately to avoid sanitization errors
umangyadav Nov 16, 2023
1cf87ef
use std::min to avoid sanitization errors
umangyadav Nov 16, 2023
e7e5ba2
Merge branch 'develop' into ref_fp8
umangyadav Nov 16, 2023
98a838f
formatting
umangyadav Nov 16, 2023
61e4e1d
use 31 for min value
umangyadav Nov 16, 2023
a5c38eb
add note
umangyadav Nov 16, 2023
61775ea
Merge branch 'ref_fp8' of github.com:ROCmSoftwarePlatform/AMDMIGraphX…
umangyadav Nov 16, 2023
3806427
Merge branch 'develop' into ref_fp8
umangyadav Nov 16, 2023
017d67e
add some more comments
umangyadav Nov 17, 2023
9e6d866
Merge branch 'ref_fp8' of github.com:ROCmSoftwarePlatform/AMDMIGraphX…
umangyadav Nov 17, 2023
a9dd42f
port gpu changes
umangyadav Nov 17, 2023
d7339e8
use bit cast
umangyadav Nov 17, 2023
6094234
Make FNUZ template param and add numeric limits
umangyadav Nov 17, 2023
78ec77e
only compile for device
umangyadav Nov 17, 2023
3411649
remove non-JIT related code
umangyadav Nov 17, 2023
d2c25a0
Remove FP8_Lowest/Max
umangyadav Nov 17, 2023
5da68df
remove using for dtypes
umangyadav Nov 17, 2023
b36f72d
Update float8_impl
umangyadav Nov 17, 2023
85ba819
constructor from float works with constexpr
umangyadav Nov 17, 2023
aed1922
Remove unnecessary pragmas
umangyadav Nov 17, 2023
f975c63
Remove clang diagnostics
umangyadav Nov 17, 2023
32033d8
Add back floatequal
umangyadav Nov 17, 2023
e88d46a
disable DPP For FP8
umangyadav Nov 17, 2023
3ae93ca
Merge remote-tracking branch 'origin/develop' into gpu_fp8
umangyadav Nov 17, 2023
60dd1f4
formatting
umangyadav Nov 17, 2023
ef425d0
revert unwanted changes
umangyadav Nov 17, 2023
76f0318
Merge branch 'gpu_fp8' of https://github.com/ROCmSoftwarePlatform/AMD…
umangyadav Nov 17, 2023
bd0ae5f
add some more tests
umangyadav Nov 17, 2023
91cc9c7
Add math and reduce tests
umangyadav Nov 18, 2023
e2b0c40
Fix tidy and other errors
umangyadav Nov 18, 2023
9f50051
fixes
umangyadav Nov 18, 2023
249464c
add nolint
umangyadav Nov 18, 2023
1be9587
tidy fix
umangyadav Nov 18, 2023
13403ab
roialign, softmax, pow, acosh, atanh,pad tests are enabled now
umangyadav Nov 20, 2023
f550f81
add layernorm, remove constexpr for 1/r
umangyadav Nov 20, 2023
7e3444c
tidy fixes
umangyadav Nov 20, 2023
6155c78
use __builtin_is_constant_evaluated
umangyadav Nov 20, 2023
13ef414
add test for rsqrt and remove old-styple-cast
umangyadav Nov 20, 2023
8660572
add comment about c++20 extensions
umangyadav Nov 20, 2023
6fbd997
Remove old cast
umangyadav Nov 20, 2023
2acd265
Remove DPP
umangyadav Nov 20, 2023
836e201
Remove MIN max overloads
umangyadav Nov 20, 2023
f9542d5
Put numeric_max and numeeric lowest into float8
umangyadav Nov 20, 2023
480288f
use void for highest to match template candidates
umangyadav Nov 21, 2023
a6c5772
add float8 for tensorview
umangyadav Nov 21, 2023
3aa465f
compiles all right
umangyadav Nov 26, 2023
037205c
Works now
umangyadav Nov 26, 2023
87548b5
add ifdef to compile
umangyadav Nov 26, 2023
d473b80
add tests and fix cmake
umangyadav Nov 26, 2023
4604f2e
add tests
umangyadav Nov 26, 2023
ad9c25e
add eliminate_fp8 pass
umangyadav Nov 26, 2023
8734ffa
remove convert from lowering
umangyadav Nov 26, 2023
f014fb9
Fix eliminate_fp8 pass
umangyadav Nov 26, 2023
83ce487
Move pass before optimize module
umangyadav Nov 26, 2023
9a9e964
formatting
umangyadav Nov 26, 2023
c40a39c
fix cppcheck
umangyadav Nov 26, 2023
c4cee34
Merge branch 'develop' into rocblas_fp8
umangyadav Dec 1, 2023
f155b0e
merge changes
umangyadav Dec 1, 2023
38218ed
few changes
umangyadav Dec 1, 2023
379692f
few more cosmetic changes
umangyadav Dec 1, 2023
381b2d9
add half tests
umangyadav Dec 2, 2023
5423577
use updated eliminate_fp8 pass
umangyadav Dec 4, 2023
402c66a
use eliminate_data_type pass instead of eliminate_fp8 pass
umangyadav Dec 5, 2023
8738f3b
Merge branch 'develop' into rocblas_fp8
umangyadav Dec 5, 2023
4ca90ec
remove older files
umangyadav Dec 5, 2023
b099a7d
remove header
umangyadav Dec 5, 2023
7d6e6ad
fix typo
umangyadav Dec 5, 2023
cf91c2b
add changes for the eliminate_data_type pass
umangyadav Dec 5, 2023
82f9847
add comments
umangyadav Dec 5, 2023
a9db2bf
fix typo
umangyadav Dec 5, 2023
aeaac20
remove else
umangyadav Dec 5, 2023
7c72d33
try merge eliminate_fp8 and eliminate_data_type
umangyadav Dec 2, 2023
ecd360d
add fill verify test, it is not doing anything though
umangyadav Dec 4, 2023
9e8930e
add device kernels
umangyadav Dec 4, 2023
83004a1
add RNN
umangyadav Dec 4, 2023
079e47f
add all device ops
umangyadav Dec 4, 2023
2ceadcd
remove fill test
umangyadav Dec 4, 2023
faf64ac
merge fixes
umangyadav Dec 5, 2023
6d479f8
remove float8 header
umangyadav Dec 5, 2023
ca47173
add back fp8 converts
umangyadav Dec 5, 2023
49edf5b
changes
umangyadav Dec 5, 2023
66db3c2
undo formatting changes
umangyadav Dec 5, 2023
ba86b0f
Merge branch 'develop' into device_fp8
umangyadav Dec 6, 2023
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
Next Next commit
use std::min to avoid sanitization errors
umangyadav committed Nov 16, 2023

Verified

This commit was created on GitHub.com and signed with GitHub’s verified signature.
commit 1cf87efbd895623cac5e9b6a45d8b15ddaa3ba91
24 changes: 10 additions & 14 deletions src/include/migraphx/float8_impl.hpp
Original file line number Diff line number Diff line change
@@ -22,6 +22,8 @@

#ifndef MIGRAPHX_GUARD_RTGLIB_FLOAT8_IMPL_HPP
#define MIGRAPHX_GUARD_RTGLIB_FLOAT8_IMPL_HPP
#include <algorithm>
#include <cstdint>
#include <type_traits>
#include <migraphx/config.hpp>
#include <migraphx/bit_cast.hpp>
@@ -149,19 +151,13 @@ constexpr uint8_t cast_to_f8(T f_x, bool stoch = false, uint32_t rng = 0)
else
{ // fp32/fp16 is normal with implicit 1
act_exponent = exponent - bias;
/*
check if FP8 is underflowing to 0.0. Wm is added to check to allow FP8 to go into denorm
range. e.g. act_exponent for FP32/16 is -9 and e4m3fnuz has denorm_act exponent = -7 in
that case fp32/16 mantissa can be shifted right by two to make
exponent -7 and then it can be representable as e4m3fnuz denorm. So for fp32/fp16, exponent
-10 is the cut point to convert to e4m3fp8fnuz due to implicit 1 in mantissa. If fp32/16
act_exponent is less than -10 then it underflows to zero*/
if(act_exponent < (f8_denormal_act_exponent - Wm))
{
return NegativeZeroNan ? 0x00 : ((sign) ? 0x80 : 0x00);
}
else if(act_exponent <= f8_denormal_act_exponent)
if(act_exponent <= f8_denormal_act_exponent)
{
/* This is the case where fp32/fp16 is normal but it is in f8 denormal range.
For example fp8 FNUZ mode, denormal exponent is -7, but if the fp32/fp16
actual exponent is -7, it is actually larger due to the implict 1,
Therefore it needs to be adjust to -6 and mantissa shift right by 1.
So for fp32/fp16, exponent -8 is the cut point to convert to fp8 FNUZ */
exponent_diff = f8_denormal_act_exponent - act_exponent;
}
else
@@ -173,8 +169,8 @@ constexpr uint8_t cast_to_f8(T f_x, bool stoch = false, uint32_t rng = 0)
mantissa += (1u << mfmt); // Add the implicit 1 into mantissa
}

bool midpoint = (mantissa & ((1 << (mfmt - Wm + exponent_diff)) - 1)) ==
(1 << (mfmt - Wm + exponent_diff - 1));
bool midpoint = (mantissa & ((1u << std::min(32u, mfmt - Wm + exponent_diff)) - 1)) ==
(1u << std::min(32u, mfmt - Wm + exponent_diff - 1));
/* This part is a bit tricky. The judgment of whether it is a tie needs to be done before we
shift right as shift right could rip off some residual part and make something not midpoint look
like midpoint. For example, the fp16 number 0x1002 (0 00100 0000000010), it is larger than