Skip to content

fix(runtime): defer MIOpen handle fetch for vendor-BLAS-disabled elem… #619

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
amd-mah wants to merge 3 commits into
base: main
Choose a base branch
from
+13 −7
Open

fix(runtime): defer MIOpen handle fetch for vendor-BLAS-disabled elem… #619

Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
522dc22
fix(runtime): defer MIOpen handle fetch for vendor-BLAS-disabled elem…
amd-mah Jul 31, 2026
51d6ce1
Merge branch 'main' into fix/vendor-blas-handle-subgraph
amd-mah Aug 3, 2026
ab74bed
Merge branch 'main' into fix/vendor-blas-handle-subgraph
amd-mah Aug 5, 2026
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
20 changes: 13 additions & 7 deletions lib/Runtime/real/elementwise.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -380,13 +380,6 @@ int wrap_miopenOpTensor(RuntimeState *state, int op_state_slot, void *lhs,
(long long)out_n, (long long)out_c, (long long)out_h,
(long long)out_w, type_name, (long long)data_type);

miopenHandle_t handle =
static_cast<miopenHandle_t>(hipdnn_ep_state_get_miopen_handle(state));
if (!handle) {
fprintf(stderr, "wrap_miopenOpTensor: null MIOpen handle\n");
return -1;
}

bool op_ok;
miopenTensorOp_t miopen_op = hipdnn_ep_to_miopen_op(tensor_op, op_ok);
if (!op_ok) {
Expand Down Expand Up @@ -655,6 +648,19 @@ int wrap_miopenOpTensor(RuntimeState *state, int op_state_slot, void *lhs,

float alpha1 = 1.0f, alpha2 = 1.0f, beta = 0.0f;

// The handle is fetched here — after the integer and float/half fast paths
// above have had their chance to return — because only the MIOpen fallback
// below actually needs it. Builds that disable the vendor backends
// (HIPDNN_EP_DISABLE_VENDOR_BLAS) leave the handle null by design and expect
// ADD/MUL/MIN/MAX to be served entirely by the custom kernels; checking the
// handle earlier would abort those ops even though they never touch MIOpen.
miopenHandle_t handle =

@fhanuman fhanuman Aug 4, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for making this change. Does it make sense to add a test to lock-down this code path?

static_cast<miopenHandle_t>(hipdnn_ep_state_get_miopen_handle(state));
if (!handle) {
fprintf(stderr, "wrap_miopenOpTensor: null MIOpen handle\n");
return -1;
}

RUNTIME_DEBUG_LOG("[REAL] wrap_miopenOpTensor: calling miopenOpTensor"
"(op=%s, alpha1=%.1f, alpha2=%.1f, beta=%.1f)\n",
op_name, alpha1, alpha2, beta);
Expand Down
Loading