diff --git a/onnxruntime/core/providers/rocm/tunable/gemm_hipblaslt.h b/onnxruntime/core/providers/rocm/tunable/gemm_hipblaslt.h index 7d9beefaad..d5f9de26ad 100644 --- a/onnxruntime/core/providers/rocm/tunable/gemm_hipblaslt.h +++ b/onnxruntime/core/providers/rocm/tunable/gemm_hipblaslt.h @@ -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>> 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 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(), "HipBlasLt_", i); + std::string type_string = onnxruntime::MakeString( + TypeStringFor(), "HipBlasLt_", i, "_algo_", algo_index); ret.emplace_back(type_string, std::move(hipblaslt_gemm_op)); } return ret; diff --git a/onnxruntime/core/providers/rocm/tunable/gemm_rocblas.h b/onnxruntime/core/providers/rocm/tunable/gemm_rocblas.h index 5d15a8a706..8e894e63c5 100644 --- a/onnxruntime/core/providers/rocm/tunable/gemm_rocblas.h +++ b/onnxruntime/core/providers/rocm/tunable/gemm_rocblas.h @@ -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>>> ret; - for (auto solution : solutions) { + for (size_t i = 0; i < solutions.size(); ++i) { + auto solution = solutions[i]; auto rocblas_gemm_op = [=](const GemmParams* 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>>> ret; - for (auto solution : solutions) { + for (size_t i = 0; i < solutions.size(); ++i) { + auto solution = solutions[i]; auto rocblas_gemm_op = [=](const BatchedGemmParams* 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>>> ret; - for (auto solution : solutions) { + for (size_t i = 0; i < solutions.size(); ++i) { + auto solution = solutions[i]; auto rocblas_gemm_op = [=](const StridedBatchedGemmParams* 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; }