Skip to content

Support <F8,B,S>, <F8B8,B,S> with BiasB for type TN #522

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
KKyang merged 4 commits into from
Jan 4, 2024
Merged

Support <F8,B,S>, <F8B8,B,S> with BiasB for type TN #522

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
2beae93
support <F8,B,S> and <F8B8,B,S> with bias B16
solaslin Dec 20, 2023
e7f30ec
added tox tests and gtests
solaslin Dec 28, 2023
642f722
update gfx941
solaslin Dec 29, 2023
f68fe6c
updated as per review comments
solaslin Jan 3, 2024
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
2 changes: 1 addition & 1 deletion CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -129,7 +129,7 @@ if(NOT BUILD_CUDA)
message(STATUS "Using hip-clang to build for amdgpu backend")
set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -D__HIP_HCC_COMPAT_MODE__=1" )
if (CMAKE_BUILD_TYPE MATCHES "Debug" AND NOT WIN32)
set (CMAKE_CXX_FLAGS_DEBUG "-O1 ${CMAKE_CXX_FLAGS_DEBUG} -gsplit-dwarf -ggdb" )
set (CMAKE_CXX_FLAGS_DEBUG "-O0 ${CMAKE_CXX_FLAGS_DEBUG} -gsplit-dwarf -ggdb" )
endif()
Comment on lines 131 to 133

@solaslin solaslin Dec 28, 2023

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

-O0 can get more detailed info when debugging

else()
message(FATAL_ERROR "'hipcc' compiler required to compile for ROCm software.")
Expand Down
4 changes: 2 additions & 2 deletions clients/benchmarks/client.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -472,11 +472,11 @@ try

("amaxScaleA",
bool_switch(&arg.amaxScaleA)->default_value(false),
"Apple scale for A buffer by abs max of A buffer")
"Apply scale for A buffer by abs max of A buffer")

("amaxScaleB",
bool_switch(&arg.amaxScaleB)->default_value(false),
"Apple scale for B buffer by abs max of B buffer")
"Apply scale for B buffer by abs max of B buffer")

("use_e",
bool_switch(&arg.use_e)->default_value(false),
Expand Down
2 changes: 2 additions & 0 deletions clients/common/cblas_interface.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -283,8 +283,10 @@ CREATEFUNCTION(hipblaslt_bf8, hipblaslt_f8, float, float, hipblaslt_bf8)
CREATEFUNCTION(hipblaslt_f8, hipblaslt_bf8, float, float, hipblaslt_f8)
CREATEFUNCTION(hipblaslt_f8, hipblaslt_f8, hipblasLtHalf, float, hipblaslt_f8)
CREATEFUNCTION(hipblaslt_f8, hipblaslt_f8, hipblaslt_f8, float, hipblaslt_f8)
CREATEFUNCTION(hipblaslt_f8, hipblaslt_f8, hip_bfloat16, float, hipblaslt_f8)
CREATEFUNCTION(hipblaslt_bf8, hipblaslt_f8, hipblasLtHalf, float, hipblaslt_bf8)
CREATEFUNCTION(hipblaslt_f8, hipblaslt_bf8, hipblasLtHalf, float, hipblaslt_f8)
CREATEFUNCTION(hipblaslt_f8, hipblaslt_bf8, hip_bfloat16, float, hipblaslt_f8)
CREATEFUNCTION(hipblasLtHalf, hipblasLtHalf, hipblasLtHalf, float, hipblasLtHalf)
CREATEFUNCTION(hipblasLtHalf, hipblasLtHalf, float, float, hipblasLtHalf)
CREATEFUNCTION(float, float, float, float, float)
Expand Down
18 changes: 18 additions & 0 deletions clients/gtest/matmul_gtest.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -601,6 +601,24 @@ Tests:
unit_check: 1
gpu_arch: '94?'

- name: matmul_f8_bf8_dst_fp16
category: pre_checkin
function:
matmul: *real_precisions_1b_dst_bf16
M: [128, 129]
N: [128, 129]
K: [128, 129]
transA: T
transB: N
alpha: 1
beta: [ 0.0, 2.0 ]
scaleA: [ 0, 1]
scaleB: [ 0, 1]
bias_vector: [0, 1]
bias_type: bf16_r
unit_check: 1
gpu_arch: '94?'

- name: matmul_f8_dst_f8
category: pre_checkin
function:
Expand Down
11 changes: 11 additions & 0 deletions clients/include/hipblaslt_common.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -72,6 +72,12 @@ Real precisions 1 bytes: &real_precisions_1b
{ a_type: bf8_r, b_type: f8_r, c_type: f16_r, d_type: f16_r, compute_type: c_f32_r, scale_type: f32_r}
- &f8bf8_precision_dst_fp16
{ a_type: f8_r, b_type: bf8_r, c_type: f16_r, d_type: f16_r, compute_type: c_f32_r, scale_type: f32_r}
- &f8_precision_dst_bf16
{ a_type: f8_r, b_type: f8_r, c_type: bf16_r, d_type: bf16_r, compute_type: c_f32_r, scale_type: f32_r}
- &f8bf8_precision_dst_bf16
{ a_type: f8_r, b_type: bf8_r, c_type: bf16_r, d_type: bf16_r, compute_type: c_f32_r, scale_type: f32_r}
- &bf8f8_precision_dst_bf16
{ a_type: bf8_r, b_type: f8_r, c_type: bf16_r, d_type: bf16_r, compute_type: c_f32_r, scale_type: f32_r}
- &f8_precision_dst_f8
{ a_type: f8_r, b_type: f8_r, c_type: f8_r, d_type: f8_r, compute_type: c_f32_r, scale_type: f32_r}

Expand All @@ -85,6 +91,11 @@ Real precisions 1 bytes dst fp16: &real_precisions_1b_dst_f16
- *bf8f8_precision_dst_fp16
- *f8bf8_precision_dst_fp16

Real precisions 1 bytes dst bf16: &real_precisions_1b_dst_bf16
- *f8_precision_dst_bf16
- *f8bf8_precision_dst_bf16
# - *bf8f8_precision_dst_bf16 # TODO

Real precisions 1 bytes dst fp8: &real_precisions_1b_dst_f8
Comment on lines 93 to 99

@solaslin solaslin Dec 28, 2023

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

the third one is not included in this pr but we are gonna support it later. so this is just a placeholder

- *f8_precision_dst_f8

Expand Down
10 changes: 10 additions & 0 deletions clients/include/type_dispatch.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -139,6 +139,11 @@ auto hipblaslt_matmul_dispatch(const Arguments& arg)
{
return TEST<hipblaslt_f8, hipblaslt_f8, hipblasLtHalf, float>{}(arg);
}
else if(TiA == HIPBLASLT_R_8F_E4M3 && TiB == HIPBLASLT_R_8F_E4M3 && To == HIPBLASLT_R_16B
&& Tc == HIPBLASLT_COMPUTE_F32)
{
return TEST<hipblaslt_f8, hipblaslt_f8, hipblasLtBfloat16, float>{}(arg);
}
else if(TiA == HIPBLASLT_R_8F_E4M3 && TiB == HIPBLASLT_R_8F_E4M3
&& To == HIPBLASLT_R_8F_E4M3 && Tc == HIPBLASLT_COMPUTE_F32)
{
Expand All @@ -154,6 +159,11 @@ auto hipblaslt_matmul_dispatch(const Arguments& arg)
{
return TEST<hipblaslt_f8, hipblaslt_bf8, hipblasLtHalf, float>{}(arg);
}
else if(TiA == HIPBLASLT_R_8F_E4M3 && TiB == HIPBLASLT_R_8F_E5M2 && To == HIPBLASLT_R_16B
&& Tc == HIPBLASLT_COMPUTE_F32)
{
return TEST<hipblaslt_f8, hipblaslt_bf8, hipblasLtBfloat16, float>{}(arg);
}
/*
else if(Ti == HIPBLASLT_R_8I && To == HIPBLASLT_R_8I && Tc == HIPBLASLT_COMPUTE_I32)
{
Expand Down
6 changes: 3 additions & 3 deletions library/include/hipblaslt.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -114,7 +114,7 @@ typedef enum {
typedef enum {
HIPBLASLT_COMPUTE_F32 = 300, /**<32-bit floating-point precision.*/
HIPBLASLT_COMPUTE_F32_FAST_XF32 = 301, /**< XF32 compute for 32-bit input and output matrices */
HIPBLASLT_COMPUTE_F64 = 302, /**<32-bit floating-point precision.*/
HIPBLASLT_COMPUTE_F64 = 302, /**<64-bit floating-point precision.*/
HIPBLASLT_COMPUTE_I32 = 303,
Comment on lines 116 to 118

@solaslin solaslin Dec 28, 2023

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

A typo ?

HIPBLASLT_COMPUTE_F32_FAST_F16 = 304, /**< F16 compute for 16-bit input and 32-bit output matrices */
} hipblasLtComputeType_t;
Expand Down Expand Up @@ -916,9 +916,9 @@ hipblasStatus_t hipblasLtMatrixTransformDescGetAttribute(
size_t sizeInBytes,
size_t* sizeWritten);

/*! \ingroup library_module
/*! \ingroup library_module
* \brief Matrix layout conversion helper
* \details
* \details
* Matrix layout conversion helper (C = alpha * op(A) + beta * op(B)),
* can be used to change memory order of data or to scale and shift the values.
* @param[in] lightHandle Pointer to the allocated hipBLASLt handle for the
Expand Down
Loading