Skip to content

Fix hipblaslt-bench latency issue when api_method is set to cpp #528

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

Closed
KKyang wants to merge 1 commit into from
Closed

Fix hipblaslt-bench latency issue when api_method is set to cpp #528

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
15 changes: 14 additions & 1 deletion projects/hipblaslt/clients/include/testing_matmul.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2760,7 +2760,7 @@ void testing_matmul_with_bias(const Arguments& arg,
int addRequest = 0;
for(size_t t = 0; t < 1; t++) // C API not supported yet
{
size_t tmpWorkspaceSize = 0;
size_t tmpWorkspaceSize = 0;
tmpAlgo[j].algo.max_workspace_bytes = max_workspace_size;
if(hipblaslt_ext::matmulIsAlgoSupported(handle,
matmul[0][0],
Expand Down Expand Up @@ -2914,6 +2914,13 @@ void testing_matmul_with_bias(const Arguments& arg,
heuristicTuningIndex.clear();
for(int j = 0; j < tmpAlgo.size(); j++)
{
if(tuningVec.size() == 1 && tuningVec[0] == hipblaslt_ext::GemmTuning())
{
heuristicResult.push_back(tmpAlgo[j]);
heuristicTuningIndex.push_back(0);
workspace_size = std::max(workspace_size, tmpAlgo[j].workspaceSize);
continue;
}
for(size_t t = 0; t < tuningVec.size(); t++)
{
size_t tmpWorkspaceSize = 0;
Expand Down Expand Up @@ -3009,6 +3016,12 @@ void testing_matmul_with_bias(const Arguments& arg,
heuristicTuningIndex.clear();
for(int j = 0; j < tmpAlgo.size(); j++)
{
if(tuningVec.size() == 1 && tuningVec[0] == hipblaslt_ext::GemmTuning())
{
heuristicResult.push_back(tmpAlgo[j]);
heuristicTuningIndex.push_back(0);
continue;
}
for(size_t t = 0; t < tuningVec.size(); t++)
{
size_t tmpWorkspaceSize = 0;
Expand Down
24 changes: 14 additions & 10 deletions projects/hipblaslt/library/include/hipblaslt/hipblaslt-ext.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -91,8 +91,9 @@ namespace hipblaslt_ext
std::unique_ptr<GemmPreferenceImpl> pimpl;
};

[[deprecated("GemmPreferenceV2 is deprecated, use GemmPreference instead.")]]
typedef GemmPreference GemmPreferenceV2;
[[deprecated(
"GemmPreferenceV2 is deprecated, use GemmPreference instead.")]] typedef GemmPreference
GemmPreferenceV2;

/*! \ingroup types_module
* \brief hipblasLt extension ProblemType for gemm problems.
Expand Down Expand Up @@ -144,8 +145,9 @@ namespace hipblaslt_ext
HIPBLASLT_EXPORT hipblasComputeType_t getTypeCompute() const; //!< The compute datatype.
};

[[deprecated("GemmProblemTypeV2 is deprecated, use GemmProblemType instead.")]]
typedef GemmProblemType GemmProblemTypeV2;
[[deprecated(
"GemmProblemTypeV2 is deprecated, use GemmProblemType instead.")]] typedef GemmProblemType
GemmProblemTypeV2;

/*! \ingroup types_module
* \brief hipblasLt extension Epilogue for gemm problems.
Expand Down Expand Up @@ -204,8 +206,8 @@ namespace hipblaslt_ext
std::unique_ptr<GemmEpilogueImpl> pimpl;
};

[[deprecated("GemmEpilogueV2 is deprecated, use GemmEpilogue instead.")]]
typedef GemmEpilogue GemmEpilogueV2;
[[deprecated("GemmEpilogueV2 is deprecated, use GemmEpilogue instead.")]] typedef GemmEpilogue
GemmEpilogueV2;

struct GemmTuning
{
Expand All @@ -215,6 +217,8 @@ namespace hipblaslt_ext

HIPBLASLT_EXPORT GemmTuning(const GemmTuning& tuning);
HIPBLASLT_EXPORT GemmTuning& operator=(const GemmTuning& tuning);
HIPBLASLT_EXPORT bool
operator==(const GemmTuning& other) const; //!< Compare two tuning objects for equality.

HIPBLASLT_EXPORT GemmTuning(GemmTuning&& tuning);
HIPBLASLT_EXPORT GemmTuning& operator=(GemmTuning&& tuning);
Expand All @@ -234,8 +238,8 @@ namespace hipblaslt_ext
std::unique_ptr<GemmTuningImpl> pimpl;
};

[[deprecated("GemmTuningV2 is deprecated, use GemmTuning instead.")]]
typedef GemmTuning GemmTuningV2;
[[deprecated(
"GemmTuningV2 is deprecated, use GemmTuning instead.")]] typedef GemmTuning GemmTuningV2;

/*! \ingroup types_module
* \brief hipblasLt extension Inputs for gemm problems.
Expand Down Expand Up @@ -296,8 +300,8 @@ namespace hipblaslt_ext
std::unique_ptr<GemmInputsImpl> pimpl;
};

[[deprecated("GemmInputsV2 is deprecated, use GemmInputs instead.")]]
typedef GemmInputs GemmInputsV2;
[[deprecated(
"GemmInputsV2 is deprecated, use GemmInputs instead.")]] typedef GemmInputs GemmInputsV2;

/*! \ingroup types_module
* \brief hipblasLt extension GPU inputs for gemm problems.
Expand Down
4 changes: 4 additions & 0 deletions projects/hipblaslt/library/src/amd_detail/hipblaslt-ext.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -382,6 +382,10 @@ namespace hipblaslt_ext
*pimpl = *tuning.pimpl;
return *this;
}
bool GemmTuning::operator==(const GemmTuning& other) const
{
return pimpl->splitK == other.pimpl->splitK && pimpl->wgm == other.pimpl->wgm;
}

GemmTuning::GemmTuning(GemmTuning&& tuning) = default;
GemmTuning& GemmTuning::operator=(GemmTuning&& tuning) = default;
Expand Down