Skip to content

WIP: TF32 Performance and Refactor #2274

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

Open
carsonbrownlee wants to merge 20 commits into
base: develop
Choose a base branch
from
Open

WIP: TF32 Performance and Refactor #2274

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
20 commits
Select commit Hold shift + click to select a range
52c13c0
adding f32xemulation component
carsonbrownlee Apr 15, 2025
298d44b
tf32 in situ conversion fix
carsonbrownlee Jul 8, 2025
1e6f132
debugging k32. fix for k16 without lds preconversions
carsonbrownlee Jul 9, 2025
678acc2
Add cvt after lr tf32 emu impl
b-shi Jul 11, 2025
42e4f15
2xMIk fix for non TN
b-shi Jul 11, 2025
acd2f0a
remove print
b-shi Jul 12, 2025
5c19f6e
SourceSwap Fix
b-shi Jul 12, 2025
e7567cc
Don't interleave cvt+pack with mfmas
b-shi Jul 12, 2025
4621d42
kstride fix
b-shi Jul 12, 2025
3a8b035
2xmik fix part2
b-shi Jul 12, 2025
e3be90a
Re-enable interleave
b-shi Jul 12, 2025
e232dae
Use dot2c instead of cvt+sub
b-shi Jul 13, 2025
f4bcfb2
Remove fp32->tf32 downcast in hipblaslt client
b-shi Jul 13, 2025
95dc4f7
TF32 origami libs with 16x16x32 and optional lsu
AlexBrownAMD Jul 14, 2025
a0c1989
Merge pull request #4 from AlexBrownAMD/tf32_perf
b-shi Jul 14, 2025
e0c5fd1
Temporarily disable warning message on missing instruction latency fo…
AlexBrownAMD Jul 14, 2025
75281e7
Merge pull request #5 from AlexBrownAMD/tf32_perf
b-shi Jul 14, 2025
2768dd6
tf32 kernels updated
carsonbrownlee Jul 14, 2025
84f6188
tf32 updating library logic files
carsonbrownlee Jul 14, 2025
ede0beb
updating tf32 kernels
carsonbrownlee Jul 14, 2025
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
4 changes: 2 additions & 2 deletions clients/gtest/matmul_gtest.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1247,7 +1247,7 @@ Tests:
beta: [ 0.0, 2.0 ]
unit_check: 1
algo_method: [2]
gpu_arch: '942'
gpu_arch: '9(42 | 50)'

- name: matmul_groupedgemm_zero_n
category: pre_checkin
Expand Down Expand Up @@ -1648,7 +1648,7 @@ Tests:
alpha: 1
beta: [ 0.0, 2.0 ]
unit_check: 1
gpu_arch: '942'
gpu_arch: '9(42 | 50)'

- name: matmul_gemm_double
category: pre_checkin
Expand Down
153 changes: 124 additions & 29 deletions clients/include/allclose.hpp
View file
Open in desktop

Large diffs are not rendered by default.

23 changes: 20 additions & 3 deletions clients/include/argument_model.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -72,7 +72,11 @@ class ArgumentModel
double cpu_us,
double norm,
double atol,
double rtol)
double rtol,
double absErrorMax,
double absErrorAvg,
double relErrorMax,
double relErrorAvg)
{
// requires enablement for frequency logging
ArgumentModel_log_performance(name_line, val_line);
Expand Down Expand Up @@ -154,6 +158,11 @@ class ArgumentModel
else
val_line << "," << rtol;
}
if (absErrorMax != ArgumentLogging::NA_value)
{
name_line << ",absErrorMax,absErrorAvg,relErrorMax,relErrorAvg";
val_line << "," << absErrorMax << "," << absErrorAvg << "," << relErrorMax << "," << relErrorAvg;
}
}
}
}
Expand All @@ -176,7 +185,11 @@ class ArgumentModel
double cpu_us = ArgumentLogging::NA_value,
double norm = ArgumentLogging::NA_value,
double atol = ArgumentLogging::NA_value,
double rtol = ArgumentLogging::NA_value)
double rtol = ArgumentLogging::NA_value,
double absErrorMax = ArgumentLogging::NA_value,
double absErrorAvg = ArgumentLogging::NA_value,
double relErrorMax = ArgumentLogging::NA_value,
double relErrorAvg = ArgumentLogging::NA_value)
{
hipblaslt_internal_ostream name_list;
hipblaslt_internal_ostream value_list;
Expand Down Expand Up @@ -261,7 +274,11 @@ class ArgumentModel
cpu_us,
norm,
atol,
rtol);
rtol,
absErrorMax,
absErrorAvg,
relErrorMax,
relErrorAvg);

if(archName != "")
{
Expand Down
49 changes: 40 additions & 9 deletions clients/include/testing_matmul.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -831,7 +831,11 @@ void check(hipStream_t stream,
hipDataType To,
hipDataType Tbias,
hipDataType Taux,
hipDataType Tc)
hipDataType Tc,
double& absErrorMax,
double& absErrorAvg,
double& relErrorMax,
double& relErrorAvg)
{
// fetch GPU
CHECK_HIP_ERROR(hipStreamSynchronize(stream));
Expand Down Expand Up @@ -1040,7 +1044,11 @@ void check(hipStream_t stream,
num_batches[gemmIdx],
hipblaslt_atol,
hipblaslt_rtol,
To);
To,
absErrorMax,
absErrorAvg,
relErrorMax,
relErrorAvg);
//TODO: confirm if allclose_check_assert is neccessary
}
}
Expand Down Expand Up @@ -3080,15 +3088,17 @@ void testing_matmul_with_bias(const Arguments& arg,
}

// For the xf32 xdl math op, cast type of A/B from float to xfloat32 .
if(TiA == HIP_R_32F && TiB == HIP_R_32F && To == HIP_R_32F && Talpha == HIP_R_32F)
if(arg.compute_type == HIPBLAS_COMPUTE_32F_FAST_TF32)
// NOTE: This should not be needed. Converting inputs to TF32 would result in
// less accurate reference values.
/*if(TiA == HIP_R_32F && TiB == HIP_R_32F && To == HIP_R_32F && Talpha == HIP_R_32F)
if(arg.compute_type == HIPBLAS_COMPUTE_32F_FAST_TF32 && false)
for(int i = 0; i < gemm_count; i++)
{
type_to_xdl_math_op_type<hipblasLtXfloat32, float, float>(
static_cast<float*>(hA[i].buf()), size_A[i]);
type_to_xdl_math_op_type<hipblasLtXfloat32, float, float>(
static_cast<float*>(hB[i].buf()), size_B[i]);
}
}*/

#define epilogue_param \
M[gemmIdx], N[gemmIdx], ldd[gemmIdx], \
Expand Down Expand Up @@ -3360,6 +3370,11 @@ void testing_matmul_with_bias(const Arguments& arg,
}
}

double absErrorMax = 0;
double absErrorAvg = 0;
double relErrorMax = 0;
double relErrorAvg = 0;

if(!arg.timing)
{
for(size_t sol = 0; sol < heuristicResult.size(); sol++)
Expand Down Expand Up @@ -3480,7 +3495,11 @@ void testing_matmul_with_bias(const Arguments& arg,
To,
Tbias,
Taux,
Talpha);
Talpha,
absErrorMax,
absErrorAvg,
relErrorMax,
relErrorAvg);
}
}
}
Expand Down Expand Up @@ -3883,7 +3902,11 @@ void testing_matmul_with_bias(const Arguments& arg,
To,
Tbias,
Taux,
Talpha);
Talpha,
absErrorMax,
absErrorAvg,
relErrorMax,
relErrorAvg);
}

#define argument_param \
Expand Down Expand Up @@ -3952,7 +3975,11 @@ void testing_matmul_with_bias(const Arguments& arg,
cpu_time_used,
hipblaslt_error,
hipblaslt_atol,
hipblaslt_rtol);
hipblaslt_rtol,
absErrorMax,
absErrorAvg,
relErrorMax,
relErrorAvg);
}
if(best_gpu_time > gpu_time_used)
{
Expand Down Expand Up @@ -4010,7 +4037,11 @@ void testing_matmul_with_bias(const Arguments& arg,
cpu_time_used,
best_norm,
best_atol,
best_rtol);
best_rtol,
absErrorMax,
absErrorAvg,
relErrorMax,
relErrorAvg);
}
}

Expand Down
Loading