Skip to content

Fp8 gemm for gfx940 enabled. #8

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.

Sign up for GitHub

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

Merged
wenchenvincent merged 3 commits into from
Aug 24, 2023
Merged

Fp8 gemm for gfx940 enabled. #8

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
32045db
Initial integration of rocBlas F8 GEMM without epilogues. Changed F8 …
wenchenvincent Apr 5, 2023
8a4d985
Added support for bias data type different from matrix B or D.
wenchenvincent Apr 14, 2023
ac69653
Enabled fp8 native conversion instructions for gfx940.
wenchenvincent Apr 20, 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
106 changes: 99 additions & 7 deletions transformer_engine/common/amd_detail/hip_float8.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,6 +2,11 @@
// FP8 header version 0.3, 2021/05/11

#define HIP_HOST_DEVICE __host__ __device__
#define HIP_DEVICE __device__
#define HIP_HOST __host__

#define E5M2_AMAX 57344.0
#define E4M3_AMAX 240.0

@HaiShaw HaiShaw Aug 23, 2023

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We may consider 224 for better OCP interop


namespace hip_f8_impl {

Expand Down Expand Up @@ -49,8 +54,8 @@ enum class hip_f8_rounding_mode {
// => bias = 15 for 152, 7 for 143
// => NAN/INF are represented as per IEEE conventions

static __device__ bool hip_f8_bias_mode_bit_device;
static bool hip_f8_bias_mode_bit_host;
static __device__ bool hip_f8_bias_mode_bit_device = true;
static bool hip_f8_bias_mode_bit_host = true;

static __global__ void set_hip_f8_bias_mode_bit(bool v) {
hip_f8_bias_mode_bit_device = v;
Expand All @@ -69,7 +74,6 @@ static void set_hip_f8_bias_mode_optimal() {
static inline HIP_HOST_DEVICE bool get_hip_f8_bias_mode() {
#if defined(__HIP_DEVICE_COMPILE__)
return hip_f8_bias_mode_bit_device;

#else
return hip_f8_bias_mode_bit_host;
#endif
Expand All @@ -89,7 +93,53 @@ struct hip_f8 {
}

// constructor from float
explicit HIP_HOST_DEVICE hip_f8(float v, hip_f8_rounding_mode rm=hip_f8_rounding_mode::standard, uint32_t rng=0) {
#if defined(__gfx940__)

@HaiShaw HaiShaw Aug 23, 2023

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

May need to consider gfx941, etc.

@wenchenvincent wenchenvincent Aug 24, 2023

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This was addressed in a PR #9

explicit HIP_DEVICE hip_f8(float v, hip_f8_rounding_mode rm=hip_f8_rounding_mode::standard, uint32_t rng=0) {
union {
float fval;
uint32_t i32val;
uint8_t i8val[4];
} val;
uint32_t ival = 0;
val.fval = v;

if (T == hip_f8_type::bf8) { // bf8
if ((val.i32val & 0x7F800000) != 0x7F800000) // propagate NAN/INF, no clipping
val.fval = __builtin_amdgcn_fmed3f(val.fval, E5M2_AMAX, -E5M2_AMAX);
if (rm == hip_f8_rounding_mode::standard) { // RNE rounding
ival = __builtin_amdgcn_cvt_pk_bf8_f32(
val.fval, val.fval, ival, false); // false -> WORD0
val.i32val = ival;
data = val.i8val[0];
}
else { //stochastic rounding
ival = __builtin_amdgcn_cvt_sr_bf8_f32(val.fval, rng, ival, 0); // 0 pos
val.i32val = ival;
data = val.i8val[0]; // little endian
}
}
else { // fp8
if ((val.i32val & 0x7F800000) != 0x7F800000) /// propagate NAN/INF, no clipping
val.fval = __builtin_amdgcn_fmed3f(val.fval, E4M3_AMAX, -E4M3_AMAX);
if (rm == hip_f8_rounding_mode::standard) { // RNE rounding
ival = __builtin_amdgcn_cvt_pk_fp8_f32(
val.fval, val.fval, ival, false); // false -> WORD0
val.i32val = ival;
data = val.i8val[0];
}
else { //stochastic rounding
ival = __builtin_amdgcn_cvt_sr_fp8_f32(val.fval, rng, ival, 0); // 0 pos
val.i32val = ival;
data = val.i8val[0]; // little endian
}
}
}

explicit HIP_HOST //Code host still uses SW simulated conversion on gfx940
#else // #if defined(__gfx940__)
explicit HIP_HOST_DEVICE // On architectures other than gfx940, both host and device still use SW simulated conversion
#endif // #if defined(__gfx940__)
hip_f8(float v, hip_f8_rounding_mode rm=hip_f8_rounding_mode::standard, uint32_t rng=0) {
if (T == hip_f8_type::bf8) {
if (get_hip_f8_bias_mode()) {
data = hip_f8_impl::cast_to_f8<2, 5, float, true/*negative_zero_nan*/, true/*clip*/>(v, (rm == hip_f8_rounding_mode::stochastic), rng);
Expand All @@ -106,7 +156,16 @@ struct hip_f8 {
}

// constructor from half
explicit HIP_HOST_DEVICE hip_f8(half v, hip_f8_rounding_mode rm=hip_f8_rounding_mode::standard, uint32_t rng=0) {
#if defined(__gfx940__)
explicit HIP_DEVICE hip_f8(half v, hip_f8_rounding_mode rm=hip_f8_rounding_mode::standard, uint32_t rng=0)
: hip_f8((float)v, rm, rng)

@HaiShaw HaiShaw Aug 23, 2023

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

have builtins for half ?
this seems to be slower.

@wenchenvincent wenchenvincent Aug 24, 2023

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

No native instruction for fp8 upcasting to half.

{
}
explicit HIP_HOST //Code host still uses SW simulated conversion on gfx940
#else // #if defined(__gfx940__)
explicit HIP_HOST_DEVICE // On architectures other than gfx940, both host and device still use SW simulated conversion
#endif // #if defined(__gfx940__)
hip_f8(half v, hip_f8_rounding_mode rm=hip_f8_rounding_mode::standard, uint32_t rng=0) {
if (T == hip_f8_type::bf8) {
if (get_hip_f8_bias_mode()) {
data = hip_f8_impl::cast_to_f8<2, 5, half, true/*negative_zero_nan*/, true/*clip*/>(v, (rm == hip_f8_rounding_mode::stochastic), rng);
Expand All @@ -126,7 +185,32 @@ struct hip_f8 {
explicit HIP_HOST_DEVICE hip_f8(hip_bfloat16 v, hip_f8_rounding_mode r=hip_f8_rounding_mode::standard, uint32_t rng=0);

// convert to float
explicit inline HIP_HOST_DEVICE operator float() const {
#if defined(__gfx940__)
HIP_DEVICE operator float() const {
union
{
float fval;
uint32_t i32val;
uint8_t i8val[4]; // dependent of endian
} val;

// assign 8bit data in position [7:0]
val.i32val = 0;
val.i8val[0] = data; // little endian

// upcast
if(T == hip_f8_type::bf8)
val.i32val = __builtin_amdgcn_cvt_f32_bf8(val.i32val, 0); // 0 pos
else // fp8
val.i32val = __builtin_amdgcn_cvt_f32_fp8(val.i32val, 0); // 0 pos

return val.fval;
}
explicit inline HIP_HOST //Code host still uses SW simulated conversion on gfx940
#else // #if defined(__gfx940__)
explicit inline HIP_HOST_DEVICE // On architectures other than gfx940, both host and device still use SW simulated conversion
#endif // #if defined(__gfx940__)
operator float() const {
if (T == hip_f8_type::bf8) {
if (get_hip_f8_bias_mode()) {
return hip_f8_impl::cast_from_f8<2, 5, float, true/*negative_zero_nan*/>(data);
Expand All @@ -143,7 +227,15 @@ struct hip_f8 {
}

// convert to half
explicit inline HIP_HOST_DEVICE operator half() const {
#if defined(__gfx940__)
explicit HIP_DEVICE inline operator half() const {
return __half(float(*this));

@HaiShaw HaiShaw Aug 23, 2023

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

same concern of performance if there were half to fp8/bf8 builtins

}
explicit inline HIP_HOST //Code host still uses SW simulated conversion on gfx940
#else // #if defined(__gfx940__)
explicit inline HIP_HOST_DEVICE // On architectures other than gfx940, both host and device still use SW simulated conversion
#endif // #if defined(__gfx940__)
operator half() const {
if (T == hip_f8_type::bf8) {
if (get_hip_f8_bias_mode()) {
return hip_f8_impl::cast_from_f8<2, 5, half, true/*negative_zero_nan*/>(data);
Expand Down
Loading