From 37329bf8c3a337b112fe37ced47cd7a2cb74d936 Mon Sep 17 00:00:00 2001 From: Liangxijun-1001 <65910327+Liangxijun-1001@users.noreply.github.com> Date: Tue, 5 Dec 2023 14:43:27 +0800 Subject: [PATCH] =?UTF-8?q?[BugFix]=20Fix=20the=20error=20of=20reloading?= =?UTF-8?q?=20the=20model=20library=20on=20the=20ROCm=20platform:=20"MIOpe?= =?UTF-8?q?n=20Error:=20No=20invoker=20was=20registered=20for=20convolutio?= =?UTF-8?q?n=20forward.=E2=80=9D=20(#16190)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit [ROCm BugFix] Fix the error of reloading the model library on the ROCm platform: MIOpen Error: No invoker was registered for convolution forward. Signed-off-by: Liangxijun-1001 Co-authored-by: Liangxijun-1001 --- src/runtime/contrib/miopen/conv_forward.cc | 21 +++++++++++++++++++++ 1 file changed, 21 insertions(+) diff --git a/src/runtime/contrib/miopen/conv_forward.cc b/src/runtime/contrib/miopen/conv_forward.cc index 139801feef15..e1d6eb61979f 100644 --- a/src/runtime/contrib/miopen/conv_forward.cc +++ b/src/runtime/contrib/miopen/conv_forward.cc @@ -189,8 +189,29 @@ TVM_REGISTER_GLOBAL("tvm.contrib.miopen.conv2d.forward") entry_ptr->conv_entry.data_type, y->shape[0], y->shape[1], y->shape[2], y->shape[3])); + // Set workspace + size_t workspace_size = 0; + MIOPEN_CALL(miopenConvolutionForwardGetWorkSpaceSize( + entry_ptr->handle, entry_ptr->conv_entry.filter_desc, entry_ptr->conv_entry.input_desc, + entry_ptr->conv_entry.conv_desc, entry_ptr->conv_entry.output_desc, &workspace_size)); + entry_ptr->conv_entry.UpdateWorkspace(workspace_size); + const float alpha = 1.f; const float beta = 0.f; + + const int request_algo_count = 4; + const bool exhaustive_search = true; + void* workspace = entry_ptr->conv_entry.workspace; + if (workspace_size == 0) workspace = nullptr; + int returned_algo_count = 0; + miopenConvAlgoPerf_t perfs[4]; + + MIOPEN_CALL(miopenFindConvolutionForwardAlgorithm( + entry_ptr->handle, entry_ptr->conv_entry.input_desc, x->data, + entry_ptr->conv_entry.filter_desc, w->data, entry_ptr->conv_entry.conv_desc, + entry_ptr->conv_entry.output_desc, y->data, request_algo_count, &returned_algo_count, + perfs, workspace, workspace_size, exhaustive_search)); + MIOPEN_CALL(miopenConvolutionForward( entry_ptr->handle, &alpha, entry_ptr->conv_entry.input_desc, x->data, entry_ptr->conv_entry.filter_desc, w->data, entry_ptr->conv_entry.conv_desc,