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

[amd-mah]

Summary

Defer the MIOpen handle fetch in wrap_miopenOpTensor until after the integer
and float/half fast paths, so element-wise ADD/MUL/MIN/MAX run on builds that
disable the vendor backends (HIPDNN_EP_DISABLE_VENDOR_BLAS, e.g. the
functional-model build).

Why

wrap_miopenOpTensor fetched the MIOpen handle at the top of the function and
returned -1 when it was null. Builds that disable the vendor backends leave
the handle null by design and expect ADD/MUL/MIN/MAX to be served entirely by
the custom kernels. The early null check aborted those ops even though they
never touch MIOpen, so element-wise subgraphs failed on the functional-model
build.

What

• Remove the handle fetch / null check from the top of wrap_miopenOpTensor.
• Re-fetch the handle (with the same null check) only just before the MIOpen
  fallback that actually uses it, after the integer and float/half fast paths
  have had their chance to return.

Test plan

• pre-commit run --files lib/Runtime/real/elementwise.cpp — all hooks pass
  (clang-format, license header, whitespace/EOL checks).
• <ADD: build + numeric/e2e result for a vendor-BLAS-disabled build exercising
  ADD/MUL/MIN/MAX, per CONTRIBUTING's "prove GPU execution" guidance.>

Notes for reviewers

Only the vendor-BLAS-disabled path changes; when the handle is non-null the
MIOpen fallback behaves exactly as before.

Checklist

• The change is focused, or links a design/series explaining its scope.
• Relevant tests were added or updated and the results are documented.
• User-facing or design documentation was updated when needed.
• Substantial AI assistance is disclosed, and I reviewed and understand the result.

[github-actions]

Thanks for opening a PR!

This project follows LLVM's incremental-development and AI-tool-use
guidance. See CONTRIBUTING.md
for the project workflow.

Before requesting review, please check that:

1. The change is focused. Substantial work links the relevant issue
   or design discussion.
2. The PR documents relevant test results and updates affected
   documentation.
3. If AI tools provided substantial assistance, the description
   explains what was assisted and how it was validated, and commit
   trailers identify the tool. The contributor has reviewed and
   understands the result.

Reviewers are assigned through
CODEOWNERS where ownership
is configured.

[github-actions]

L2 Accuracy Results (EP vs CPU)

Model	Combined L2	Total Elems	Skipped NaN/Inf
conv_test_hybrid	4.8668E-07	64	0
GroupQueryAttention_seq256	25.2366	2621440	0
MatMulNBits_o_seq128	259.906	368640	0
QMoE_seq128	34.957	368640	0

Threshold: 0.01 | Run: 3669 - Commit: ab74bed

[github-actions]

MorphiZen EP Performance Results

Model	QPS	Session (s)	1st Infer (ms)	CPU%	Mem (MB)
full_model_seq128	7.53	9.98	363	3	1245
GroupQueryAttention_seq128	4422.25	1.67824	11	6	311
matmul_down_seq128	527.13	2.32	68	3	352

EPContext Export Performance

Model	QPS	Session (s)	1st Infer (ms)	CPU%	Mem (MB)
full_model_seq128	7.58	43.53	357	3	15589

EPContext Import Performance

Model	QPS	Session (s)	1st Infer (ms)	CPU%	Mem (MB)
full_model_seq128	7.56	9.12	357	3	15760

OGA Benchmark Results

Model	Warmup	Reps	Prompt Len	Gen Tokens	TTFT (ms)	TPS	Peak Mem (GB)	GPU Mem (GB)
gpt-oss-20b-webgpu-int4-rtn-block-32	1	5	128	128	189.0	81.1	1.33	13.53
Llama-3.1-8B-awq-g128-int4-asym-fp16-onnx-dml	1	5	128	128	239.3	40.5	1.22	6.43

OGA Wheel Smoke (Python benchmark_e2e.py)

Model	TTFT (ms)	TPS
Llama-3.1-8B-awq-g128-int4-asym-fp16-onnx-dml	129	39.9

Run: 3669 - Commit: ab74bed

[fhanuman]

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