saymrwulf/onnxruntime
1
0
Fork 0
mirror of https://github.com/saymrwulf/onnxruntime.git synced 2026-07-25 19:48:11 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

[ROCm] Sort candidate solutions in rocBLAS/hipBLASLt for deterministic offline tuning (#17297)

Browse source 
### Description

Sort the candidates in rocBLAS/hipBLASLt to make sure that they are
properly ordered and can be correctly fetched by saved indices in
offline tuning cases.
This commit is contained in:
mindest 2023-08-28 16:34:21 +08:00 committed by GitHub
parent bf8b1681f9
commit 53169f59e5
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
2 changed files with 40 additions and 20 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

14
onnxruntime/core/providers/rocm/tunable/gemm_hipblaslt.h
View file

@ -119,10 +119,18 @@ auto GetHipBlasLtTypeStringAndOps(ActivationType activation_type = ActivationTyp
heuristic_result));
HIPBLASLT_CALL_THROW(hipblasLtDestroy(handle));
// Sort heuristic_result by algo index to make sure the order of returned algos is deterministic.
std::sort(heuristic_result.begin(),
heuristic_result.end(),
[](hipblasLtMatmulHeuristicResult_t& a, hipblasLtMatmulHeuristicResult_t& b) {
return hipblaslt_ext::getIndexFromAlgo(a.algo) < hipblaslt_ext::getIndexFromAlgo(b.algo);
});
int returned_algo_count = heuristic_result.size();
std::vector<std::pair<std::string, Op<ParamsT>>> ret;
for (int i = 0; i < returned_algo_count; i++) {
hipblasLtMatmulAlgo_t algo = heuristic_result[i].algo;
int algo_index = hipblaslt_ext::getIndexFromAlgo(algo);
auto hipblaslt_gemm_op = [=](const ParamsT* params) -> Status {
hipblasLtHandle_t op_handle;
HIPBLASLT_RETURN_IF_ERROR(hipblasLtCreate(&op_handle));
@ -212,7 +220,8 @@ auto GetHipBlasLtTypeStringAndOps(ActivationType activation_type = ActivationTyp
workspace_size);
TUNABLE_OP_RETURN_UNSUPPORTED_ARGUMENT_IF(
status != HIPBLAS_STATUS_SUCCESS, "hipBLASLt find_all: algo not supported, index ", std::to_string(i));
status != HIPBLAS_STATUS_SUCCESS,
"[hipBLASLt] Solution #", i, " failed: algo ", algo_index, " not supported (", params->Signature(), ")");
IAllocatorUniquePtr<void> workspace_buffer;
if (workspace_size > 0) {
@ -243,7 +252,8 @@ auto GetHipBlasLtTypeStringAndOps(ActivationType activation_type = ActivationTyp
HIPBLASLT_RETURN_IF_ERROR(hipblasLtDestroy(op_handle));
return Status::OK();
};
std::string type_string = onnxruntime::MakeString(TypeStringFor<T, ParamsT>(), "HipBlasLt_", i);
std::string type_string = onnxruntime::MakeString(
TypeStringFor<T, ParamsT>(), "HipBlasLt_", i, "_algo_", algo_index);
ret.emplace_back(type_string, std::move(hipblaslt_gemm_op));
}
return ret;

46
onnxruntime/core/providers/rocm/tunable/gemm_rocblas.h
View file

@ -141,8 +141,12 @@ auto GetRocBlasGemmTypeStringAndOps() {
ROCBLAS_CALL_THROW(rocblas_destroy_handle(handle));
// Sort the solutions in ascending order to make the solution vector deterministic across runs
std::sort(solutions.begin(), solutions.end());
std::vector<std::pair<std::string, Op<GemmParams<T>>>> ret;
for (auto solution : solutions) {
for (size_t i = 0; i < solutions.size(); ++i) {
auto solution = solutions[i];
auto rocblas_gemm_op = [=](const GemmParams<T>* params) -> Status {
auto h_a = DoCastForHalfOrBfloat16(params->alpha);
auto h_b = DoCastForHalfOrBfloat16(params->beta);
@ -163,14 +167,14 @@ auto GetRocBlasGemmTypeStringAndOps() {
rocblas_gemm_flags_none);
TUNABLE_OP_RETURN_UNSUPPORTED_ARGUMENT_IF(
status == rocblas_status_invalid_size, "Solution ", solution, " not supported: INVALID VALUE.");
TUNABLE_OP_RETURN_UNSUPPORTED_ARGUMENT_IF(
status != rocblas_status_success, "Solution ", solution, " failed: ", rocblas_status_to_string(status));
status != rocblas_status_success,
"[rocBLAS] Solution #", i, " (original ", solution, ") failed: ", rocblas_status_to_string(status),
" (", params->Signature(), ")");
return Status::OK();
};
ret.emplace_back(std::make_pair(onnxruntime::MakeString("RocBlasGemm_", solution), std::move(rocblas_gemm_op)));
ret.emplace_back(std::make_pair(
onnxruntime::MakeString("RocBlasGemm_", i, "_sol_", solution), std::move(rocblas_gemm_op)));
}
return ret;
}
@ -206,8 +210,12 @@ auto GetRocBlasBatchedGemmTypeStringAndOps() {
ROCBLAS_CALL_THROW(rocblas_destroy_handle(handle));
// Sort the solutions in ascending order to make the solution vector deterministic across runs
std::sort(solutions.begin(), solutions.end());
std::vector<std::pair<std::string, Op<BatchedGemmParams<T>>>> ret;
for (auto solution : solutions) {
for (size_t i = 0; i < solutions.size(); ++i) {
auto solution = solutions[i];
auto rocblas_gemm_op = [=](const BatchedGemmParams<T>* params) -> Status {
auto h_a = DoCastForHalfOrBfloat16(params->alpha);
auto h_b = DoCastForHalfOrBfloat16(params->beta);
@ -229,15 +237,14 @@ auto GetRocBlasBatchedGemmTypeStringAndOps() {
rocblas_gemm_flags_none);
TUNABLE_OP_RETURN_UNSUPPORTED_ARGUMENT_IF(
status == rocblas_status_invalid_size, "Solution ", solution, " not supported: INVALID VALUE.");
TUNABLE_OP_RETURN_UNSUPPORTED_ARGUMENT_IF(
status != rocblas_status_success, "Solution ", solution, " failed: ", rocblas_status_to_string(status));
status != rocblas_status_success,
"[rocBLAS] Solution #", i, " (original ", solution, ") failed: ", rocblas_status_to_string(status),
" (", params->Signature(), ")");
return Status::OK();
};
ret.emplace_back(std::make_pair(
onnxruntime::MakeString("RocBlasBatchedGemm_", solution), std::move(rocblas_gemm_op)));
onnxruntime::MakeString("RocBlasBatchedGemm_", i, "_sol_", solution), std::move(rocblas_gemm_op)));
}
return ret;
}
@ -273,8 +280,12 @@ auto GetRocBlasStridedBatchedGemmTypeStringAndOps() {
ROCBLAS_CALL_THROW(rocblas_destroy_handle(handle));
// Sort the solutions in ascending order to make the solution vector deterministic across runs
std::sort(solutions.begin(), solutions.end());
std::vector<std::pair<std::string, Op<StridedBatchedGemmParams<T>>>> ret;
for (auto solution : solutions) {
for (size_t i = 0; i < solutions.size(); ++i) {
auto solution = solutions[i];
auto rocblas_gemm_op = [=](const StridedBatchedGemmParams<T>* params) -> Status {
auto h_a = DoCastForHalfOrBfloat16(params->alpha);
auto h_b = DoCastForHalfOrBfloat16(params->beta);
@ -296,15 +307,14 @@ auto GetRocBlasStridedBatchedGemmTypeStringAndOps() {
rocblas_gemm_flags_none);
TUNABLE_OP_RETURN_UNSUPPORTED_ARGUMENT_IF(
status == rocblas_status_invalid_size, "Solution ", solution, " not supported: INVALID VALUE.");
TUNABLE_OP_RETURN_UNSUPPORTED_ARGUMENT_IF(
status != rocblas_status_success, "Solution ", solution, " failed: ", rocblas_status_to_string(status));
status != rocblas_status_success,
"[rocBLAS] Solution #", i, " (original ", solution, ") failed: ", rocblas_status_to_string(status),
" (", params->Signature(), ")");
return Status::OK();
};
ret.emplace_back(std::make_pair(
onnxruntime::MakeString("RocBlasStridedBatchedGemm_", solution), std::move(rocblas_gemm_op)));
onnxruntime::MakeString("RocBlasStridedBatchedGemm_", i, "_sol_", solution), std::move(rocblas_gemm_op)));
}
return ret;
}