fix(gfx): copy subExecParam from device before reading timing data

[fcui-amd]

Motivation

The GFX executor crashes with a SIGSEGV when collecting per-Transfer timing
data on GPUs that do not expose VRAM to the host (e.g. PCIe cards without
large/resizable BAR).

Technical Details

Copy the subExecutor parameters back to host with hipMemcpy(..., DeviceToHost)
before reading timing/CU data, in both affected paths. No behavior change on
hardware where direct access previously worked.

Test Plan

Test Result

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[Copilot]

Pull request overview

This PR fixes a crash in the GPU GFX executor when collecting per-Transfer/per-iteration timing and CU data on systems where VRAM is not host-accessible (e.g., no large/resizable BAR) by copying SubExecParam back from device before reading fields on the host.

Changes:

• Copy rss.subExecParamGpuPtr from device to host before collecting per-iteration CU IDs in multi-stream mode.
• Copy exeInfo.subExecParamGpu from device to host before computing per-Transfer timing/CU IDs in single-stream mode.

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

[fcui-amd]

ping?

[gilbertlee-amd]

This was somewhat intentional, to try to avoid explicit SDMA traffic during the main loop.

The solution for non-AMD hardware was actually to use [managed memory instead] for the GPU parameter buffer.

TransferBench/src/header/TransferBench.hpp

Line 4371 in 2bc42cd

MemType memType = MEM_MANAGED; // NVIDIA hardware requires managed memory to access from host

I think this PR needs to be more targeted, and only be applied to GPUs that require this functionality.
Does isLargeBar report 0 for your device?

[fcui-amd]

It's amd gpu with large bar disabled.