Skip to content

Fix data transfer routing to the selected backend - #80

Open
aditya-dl wants to merge 3 commits into
onnxruntime:mainfrom
aditya-dl:fix/lazy-data-transfer-resolution
Open

Fix data transfer routing to the selected backend#80
aditya-dl wants to merge 3 commits into
onnxruntime:mainfrom
aditya-dl:fix/lazy-data-transfer-resolution

Conversation

@aditya-dl

@aditya-dl aditya-dl commented Aug 11, 2026

Copy link
Copy Markdown
Collaborator

Two shipped builds are broken by the same field:

  • 9D — Gigapixel and PhotoAI models that fall back to DirectML error out with HIP failure 1: invalid argument whenever another model in the process used MIGraphX first. Reported as AIRADSW-750, seen on RC3 and RC5.
  • 10D — no model loads at all through OGA. Every environment-level copy fails with "Data transfer implementation between source and destination device was not found."

Both are DataTransfer serving the wrong backend, or none.

ORT calls CreateDataTransfer twice: once per factory at library registration, before any CreateEp, and once per session after it. Only the registration-time instance is reachable from OrtApi::CopyTensors.

Build Backend resolution Result
9D lazy, latched on first use first backend wins forever -> DirectML copies go to HIP
10D eager, in the ctor registration-time instance freezes a null -> every env-level copy fails
this PR hybrid: freeze per-session, lazy at registration each session keeps its own backend; env-level copies still work

The two instances need opposite behaviour, so the ctor branches on whether a backend exists yet. Freezing is safe per-session because nothing can later make the answer wrong — session 2 creating a backend doesn't change which one session 1 belongs to.

Same shape as Allocator::GetBackendAllocator(), which 3fb9104 left lazy — which is why allocation succeeded on 10D and only the copy failed.

Also fixes allocators, which had the same problem: gpu_ep::Allocator re-read the shared slot on every Alloc/Free, so a second session's CreateEp redirected the first session's allocations. Now wired through OrtEp::CreateAllocator for every profile (previously DirectML only), which ORT prefers and which receives the OrtEp pointer, so each session resolves against its own backend.

Validated: two sessions constructed on separate threads, sessions created on the main thread and Run from workers, sequential two-session repro, ModelBench (PhotoAI, Gigapixel), OGA.

Known limitation: env-level copies in a mixed-backend process can still be mis-routed. No backend exposes tensor ownership, and both CanCopy return true for any AMD GPU. Not exercised by ModelBench or OGA (one model per process).

Fixes AIRADSW-750 and the 10D "Data transfer implementation between source and
destination device was not found" failure. Both are the same field read at the
wrong time.

ORT calls OrtEpFactory::CreateDataTransfer at two different times: once per
factory at library registration (Environment::RegisterExecutionProviderLibrary),
before any CreateEp has run, and once per session after CreateEp. Only the
registration-time instance is reachable from the env-level OrtApi::CopyTensors.

Three behaviours were in circulation:

  pre-3fb9104   lazy but latch-once: `if (backend_data_transfer_ == nullptr)`,
                no pointer comparison. The first use latches a backend forever,
                so a DirectML model in a process that already ran MIGraphX gets
                the HIP transfer and hipMemcpy fails with "invalid argument" on
                a DML allocation handle. This is AIRADSW-750.

  3fb9104       eager snapshot in the ctor. Fixes 750 by freezing each
                per-session transfer to its own session's backend, but the
                registration-time instance is constructed before any backend
                exists, so it freezes a null and CanCopy returns false forever.
                Every env-level copy then fails; OGA cannot load any model.

  this change   lazy plus compare-and-recreate. Re-query
                factory_.GetBackendFactory() on each use and rebuild the backend
                transfer when the pointer changes. Neither the null nor a stale
                backend can be latched.

Mirrors Allocator::GetBackendAllocator(), which 3fb9104 left lazy -- the reason
allocation succeeded and only the copy failed on 10D.

Validated on gfx1201 with ModelBench -m AllGigapixel -e amdgpu --enable-cpu:
9D RC5 errors at shared/hip/data_transfer.cc:73 on every DirectML model, this
build passes all 8 (2 MIGraphX, 6 DirectX). OGA LLM decode unaffected
(427.81 tk/s vs 413.23/398.64 baseline).

Known limitation: ModelBench runs sessions sequentially, so the shared
ProviderFactory::backend_ep_factory_ slot always matches the running model. Two
models on different backends alive and copying alternately would still
mis-route, because neither backend's CanCopy can detect a foreign allocation
(both test only device type and vendor id). Closing that needs per-session
backend state or allocation-ownership information no backend exposes today.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

@Zhaeong Zhaeong left a comment

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.

Concurrent case will break, see code for that test case here:
https://github.amd.com/3D-Perf-Tools/ModelInferencingScripts/blob/main/onnx/inferences_resnet50_onnx.cpp
https://github.amd.com/3D-Perf-Tools/ModelInferencingScripts/blob/main/onnx/inferences_resnet50_onnx_concurrent.cpp

  1. Session 1 CreateEp → field = DML.
  2. Session 2 CreateEp → field = MGX (overwrites).
  3. Session 1 Run → the shared DataTransfer sees MGX, routes the DML-allocated device
    pointer into MIGraphX's HIP hipMemcpy → invalid argument.

The lazy resolution added in the previous commit fixed env-level copies
(10D, AIRADSW-750) but regressed the two-session case: ProviderFactory::
backend_ep_factory_ is a single process-global slot, overwritten by every
session's CreateEp, so re-resolving on each use made a per-session
DataTransfer follow whichever session selected a backend most recently.
With a DirectX and a MIGraphX session alive at once, the DirectX session's
Run() passed a DML allocation handle to hipMemcpy and died with
"HIP failure 1: invalid argument" at shared/hip/data_transfer.cc:73.

ORT constructs DataTransfer at two times that need opposite behaviour:
per-session (after its own CreateEp, backend already selected) and at
library registration (before any CreateEp, no backend yet). Branch on
whether a backend exists at construction: snapshot and freeze if so,
stay lazy if not. The per-session instance's backend cannot change, so
freezing loses nothing; the registration-time instance still resolves
once a backend appears.

Allocator keeps lazy resolution deliberately -- allocators are
factory-owned and shared across sessions, so there is no single session
to pin to. Comment added there to record the asymmetry.

Verified on RX 9070 XT / gfx1201, all three arms against this build
(amdgpu-ep.dll md5 d27930ec..., migraphx-backend.dll md5 fdcb3075...):
- Two-session repro: both sessions run, no line-73 failure, both predict
  class 446 @ logit 7.16016.
- ModelBench AllGigapixel: passing (2026-08-12) -- AIRADSW-750 stays fixed.
- OGA 10D arm: og.Generator(...) constructs and decodes; the env-level
  CopyTensors path that 10D broke is working.

Two residuals, both narrower than the status quo: env-level copies in a
mixed-backend process remain mis-routable (no backend exposes tensor
ownership), and concurrent session *construction* is still racy, so the
honest claim is correctness for sequentially-constructed sessions.
@aditya-dl

Copy link
Copy Markdown
Collaborator Author

@Zhaeong fixed

@aditya-dl
aditya-dl requested a review from Zhaeong August 12, 2026 23:34
@aditya-dl

aditya-dl commented Aug 12, 2026

Copy link
Copy Markdown
Collaborator Author

@Zhaeong walking your exact sequence through 196cccc:

  1. Session 1 CreateEp → field = DML. Session 1's DataTransfer is constructed after this, sees a backend already selected, so it snapshots DML and freezes.
  2. Session 2 CreateEp → field = MGX (still overwrites — the shared slot is unchanged).
  3. Session 1 RunGetBackendDataTransfer() returns the frozen DML snapshot without consulting the shared field. DML pointer goes to DML.

The freeze is safe because session 1's backend can't change after its own CreateEp — step 2 doesn't make the earlier answer wrong. The registration-time instance stays lazy, since it's constructed before any CreateEp and is the only one OrtApi::CopyTensors can reach.

Verified with the concurrent repro you linked: both sessions run, no failure at data_transfer.cc:73. Session 1's outputs match the eager 10D build to fp rounding (-1.3252 vs -1.32422), so routing is restored rather than the error just suppressed. ModelBench (PhotoAI, Gigapixel) and the OGA arm also still pass.

@Zhaeong

Zhaeong commented Aug 13, 2026

Copy link
Copy Markdown
Collaborator

@Zhaeong walking your exact sequence through 196cccc:

  1. Session 1 CreateEp → field = DML. Session 1's DataTransfer is constructed after this, sees a backend already selected, so it snapshots DML and freezes.
  2. Session 2 CreateEp → field = MGX (still overwrites — the shared slot is unchanged).
  3. Session 1 RunGetBackendDataTransfer() returns the frozen DML snapshot without consulting the shared field. DML pointer goes to DML.

The freeze is safe because session 1's backend can't change after its own CreateEp — step 2 doesn't make the earlier answer wrong. The registration-time instance stays lazy, since it's constructed before any CreateEp and is the only one OrtApi::CopyTensors can reach.

Verified with the concurrent repro you linked: both sessions run, no failure at data_transfer.cc:73. Session 1's outputs match the eager 10D build to fp rounding (-1.3252 vs -1.32422), so routing is restored rather than the error just suppressed. ModelBench (PhotoAI, Gigapixel) and the OGA arm also still pass.

Hey @aditya-dl, it works for the session creation reverse order case now, where directx session is created, then migraphx session, then directx session runs inference, etc. Note: I also had to include your change in this PR for it to show results for both sessions:
#86

However for the concurrent case given by this example here, I just updated:
https://github.amd.com/3D-Perf-Tools/ModelInferencingScripts/blob/main/onnx/inferences_resnet50_onnx_concurrent.cpp

It still fails for me:
Running inference session 1 (profile=directx)...
[session 1] Input name: pixel_values, Output name: logits
[session 1] Output shape: 1 1000
[session 1] Predicted class index: 905 (logit=-0.0915527)
[session 1] First 10 output values: -6.80078 -7.17578 -8.41406 -8.20312 -7.98438 -6.56641 -9.09375 -7.87891 -8.46875 -9.41406
Running inference session 2 (profile=migraphx)...
[session 2] ONNX Runtime error: HIP failure 1: invalid argument; GPU=0; file=C:\Users\owen_\Documents\GitHub\onnxruntime-ep-amdgpu\src\shared\hip\data_transfer.cc; line=73; expr=hipMemcpy(dst_data, src_data, bytes, hipMemcpyHostToDevice);

Freezing the per-session DataTransfer (previous commit) fixed transfers, but
allocators still followed the process-global ProviderFactory::backend_ep_factory_
slot. gpu_ep::Allocator::GetBackendAllocator re-reads that slot on every
Alloc/Free, so with two profiles alive the second session's CreateEp redirected
the first session's allocations to the wrong backend. Reported on the PR with a
repro that constructs both sessions on their own threads: the migraphx session
died with "HIP failure 1: invalid argument" at shared/hip/data_transfer.cc:73.

ORT prefers OrtEp::CreateAllocator over OrtEpFactory::CreateAllocator when the
former is set (plugin_ep/ep_plugin_provider_interfaces.cc), and the OrtEp variant
receives the OrtEp pointer. That is enough to resolve per session: this EP already
captures its own backend at CreateEp time (backend_ep_, backend_ep_factory_).

So wire OrtEp::CreateAllocator unconditionally rather than only for DirectML, and
have ExecutionProvider::CreateAllocator try the backend's own OrtEp hook first
(DirectML implements one, for its per-session DmlBucketizedBufferAllocator) and
fall back to the backend factory (migraphx and hip implement CreateAllocator
there). Both paths go through per-EP state; neither consults GetBackendFactory().

Note the explicit null check on backend_ep_->CreateAllocator: EP_CALL_S treats a
null function pointer as success, which would return STATUS_OK with *allocator
left unset.

gpu_ep::Allocator and ProviderFactory::CreateAllocator stay as they are. They
remain reachable when ORT asks the factory directly, notably the environment-level
shared allocator, where there is no session to resolve against.

Per-thread state was tried first and rejected: Alloc/Free run on whichever thread
calls Run(), so a thread_local slot reads null on a worker thread. That produced
zeroed outputs and a GPU watchdog reset (LiveKernelEvent 141) in the
create-on-main-thread, run-on-worker-thread case.

Verified on RX 9070 XT / gfx1201, with the two commits that follow this branch
point applied (per-Compute sync + D2H stream ordering), since session 2 output is
zeroed without them:
- create sessions on main thread, Run each on a worker: both correct (was zeros)
- construct both sessions on separate threads: both correct (was line-73 failure)
- sequential two-session repro: both run, no line-73
- EP library registered on a thread that creates no session: both correct
- ModelBench AllGigapixel: passes
- OGA DeepSeek-1L, 8 tokens, 5 iterations: byte-identical to ep.9D
@aditya-dl

Copy link
Copy Markdown
Collaborator Author

@Zhaeong pushed a third commit that fixes the concurrent case you hit.

Allocators had the same shared-slot problem as DataTransfer, but freezing doesn't work for them — Alloc/Free run on whatever thread calls Run(). They now resolve through OrtEp::CreateAllocator, which receives the OrtEp pointer, so each session uses its own backend regardless of thread.

Your threaded repro passes: both sessions run, no failure at line 73. Also passes with sessions created on the main thread and Run from worker threads.

Note the run-on-worker case needs #86 as well to show non-zero session 2 output — same thing you saw.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants