[AIMIGRAPHX-828] Cross Compile Pt 2: construct gpu target with gfx number to enable cross compilation by kahmed10 · Pull Request #4795 · ROCm/AMDMIGraphX

kahmed10 · 2026-04-17T04:16:22Z

Motivation

builds on top of part 1: #4771

Technical Details

adds option to construct target with gfx arch string (ex. "gfx1100"), number of CUs, number of chiplets.
when constructed this way, benchmarking kernels in compile_ops and finalize() will be disabled.
will not work for MIOpen/hipblaslt/CK.

Changelog Category

Add a CHANGELOG.md entry for any option other than Not Applicable

- Added: New functionality.
- Changed: Changes to existing functionality.
- Removed: Functionality or support that has been removed. (Compared to a previous release)
- Optimized: Component performance that has been optimized or improved.
- Resolved Issues: Known issues from a previous version that have been resolved.
- Not Applicable: This PR is not to be included in the changelog.

…vice_name_change

codecov · 2026-04-17T06:01:36Z

Codecov Report

❌ Patch coverage is 43.28358% with 38 lines in your changes missing coverage. Please review.

Files with missing lines	Patch %	Lines
src/api/api.cpp	0.00%	15 Missing ⚠️
src/register_target.cpp	0.00%	15 Missing ⚠️
src/include/migraphx/target.hpp	85.00%	3 Missing ⚠️
src/program.cpp	40.00%	3 Missing ⚠️
src/include/migraphx/register_target.hpp	0.00%	2 Missing ⚠️

Additional details and impacted files

@@             Coverage Diff             @@
##           develop    #4795      +/-   ##
===========================================
- Coverage    92.66%   92.56%   -0.10%     
===========================================
  Files          588      588              
  Lines        30397    30467      +70     
===========================================
+ Hits         28165    28200      +35     
- Misses        2232     2267      +35

Files with missing lines	Coverage Δ
src/api/include/migraphx/migraphx.hpp	`99.03% <ø> (ø)`
src/include/migraphx/context.hpp	`83.52% <100.00%> (+2.03%)`	⬆️
src/include/migraphx/register_target.hpp	`80.00% <0.00%> (-20.00%)`	⬇️
src/include/migraphx/target.hpp	`78.82% <85.00%> (+7.16%)`	⬆️
src/program.cpp	`81.50% <40.00%> (-0.19%)`	⬇️
src/api/api.cpp	`72.41% <0.00%> (-0.83%)`	⬇️
src/register_target.cpp	`64.58% <0.00%> (-29.36%)`	⬇️

🚀 New features to boost your workflow:

❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

…oss_compile_ctx

kahmed10 · 2026-04-23T20:09:23Z

@pfultz2 @CharlieL7 this is the implementation using environment variables. I'm thinking of instead adding it to compile options with an additional map that can be used to tweak number of CUs, chiplets, etc. Also let me know what you think of the rest of the changes.

Copilot

Pull request overview

Adds a GPU cross-compilation mode driven by environment variables so MIGraphX can compile for a specified gfx* target without depending on a physical device, primarily by synthesizing GPU device properties and disabling tuning/finalization steps that require execution.

Changes:

Introduces MIGRAPHX_GPU_ARCH (plus CU/chiplet overrides) and a synthetic HIP device-props path to enable GPU cross-compilation.
Disables kernel benchmarking/tuning paths when cross-compiling (problem cache load + compile_ops benchmarking).
Updates GPU passes and data-type elimination logic to query device capabilities via context where available; documents the new env vars and adds a changelog entry.

Reviewed changes

Copilot reviewed 13 out of 13 changed files in this pull request and generated 6 comments.

Show a summary per file

File	Description
src/targets/gpu/target.cpp	Adds env vars, creates cross-compile GPU context, and avoids problem-cache load in cross-compile mode.
src/targets/gpu/include/migraphx/gpu/eliminate_data_type_for_gpu.hpp	Adds `context*` plumbing to support context-based device capability queries.
src/targets/gpu/include/migraphx/gpu/cross_compile_device.hpp	Declares helper to build synthetic `hipDeviceProp_t` for cross compilation.
src/targets/gpu/include/migraphx/gpu/context.hpp	Adds cross-compile constructors/mode and blocks runtime execution paths in that mode.
src/targets/gpu/eliminate_data_type_for_gpu.cpp	Uses context-aware queries (when provided) for fp8/bf16 capability checks.
src/targets/gpu/device_name.cpp	Refactors hipBLASLt support helpers and retains context/no-context query variants.
src/targets/gpu/cross_compile_device.cpp	Implements synthetic HIP device properties for cross compilation.
src/targets/gpu/compile_ops.cpp	Skips benchmarking/tuning when cross-compiling.
src/targets/gpu/CMakeLists.txt	Adds new cross-compilation device source to the GPU library.
src/program.cpp	Skips finalize/module-finalize when `MIGRAPHX_GPU_ARCH` is set.
src/include/migraphx/program.hpp	Declares `MIGRAPHX_GPU_ARCH` env var for core program compilation logic.
docs/reference/MIGraphX-dev-env-vars.rst	Documents `MIGRAPHX_GPU_ARCH`, CU/chiplet overrides, and stated limitations.
CHANGELOG.md	Adds a changelog entry for cross-compilation support.

@@ -342,7 +344,8 @@ void program::compile(const target& t, compile_options options)
            MIGRAPHX_THROW("Dangling reference in module " + mod->name() + " from instruction " +
                           std::to_string(index));
        }
-        mod->finalize(this->impl->contexts);
+        if(not cross_compiling)
+            mod->finalize(this->impl->contexts);


+    if(not arch.empty())
+    {
+        auto num_cu       = value_of(MIGRAPHX_GPU_NUM_CU{}, 120);
+        auto num_chiplets = value_of(MIGRAPHX_GPU_NUM_CHIPLETS{}, 1);


+#include <migraphx/gpu/export.h>
+#include <migraphx/config.hpp>
+#include <hip/hip_runtime_api.h>
+#include <string>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace gpu {
+
+/// Populate a hipDeviceProp_t with synthetic values for cross-compilation.
+/// Used when no physical GPU is present and the target architecture
+/// is specified via environment variables.
+MIGRAPHX_GPU_EXPORT hipDeviceProp_t make_cross_compile_device_props(const std::string& arch_name,
+                                                                    std::size_t cu_count);


 struct MIGRAPHX_GPU_EXPORT eliminate_data_type_for_gpu
 {
    bool disable_64bit = false;
+    const context* ctx = nullptr;
    std::string name() const { return "gpu::eliminate_data_type_for_gpu"; }
    void apply(module_pass_manager& mpm) const;


bdevorem · 2026-05-28T16:04:23Z

+      | Enables cross-compilation mode by specifying a target GPU architecture without requiring a physical GPU.
+      | When set, kernel benchmarking and finalization are skipped. MIOpen, hipBLASLt, and CK operations are currently not supported in this mode.
+
+    - | Takes a valid GPU architecture string (e.g. ``gfx942``, ``gfx1100``).
+
+      | Default: Not set. A physical GPU is used.


@kahmed10 is the intent to let it throw in cross compile mode for external lib calls or should it be handled nicer?

+    if(string_value_of(MIGRAPHX_GPU_ARCH{}).empty())
+        this->finalize();


pfultz2 · 2026-04-24T20:41:28Z

    ctx.set_exhaustive_tune_flag(options.exhaustive_tune);
-    ctx.load_problem_cache();
+    if(not ctx.is_cross_compile())
+        ctx.load_problem_cache();


The problem_cache should still be loaded when cross compiling. Otherwise compile_ops will not pick the fastest config.

pfultz2 · 2026-04-24T20:46:08Z

        }
    }
-    this->finalize();
+    if(string_value_of(MIGRAPHX_GPU_ARCH{}).empty())


I dont think we should use env variable here. Instead the target or context should tell us if we can run finalize.

pfultz2 · 2026-04-24T20:47:04Z

    auto&& passes = t.get_passes(this->impl->contexts.front(), options);
    run_passes(*this, passes, options.trace);
    auto mods = this->get_modules();
+    bool cross_compiling = not string_value_of(MIGRAPHX_GPU_ARCH{}).empty();


I dont think we should use env variable here. Instead the target or context should tell us if we can run finalize.

pfultz2 · 2026-04-24T20:52:00Z

+    if(ctx != nullptr)
+        return f(*ctx);
+    return f();
+}


Remove this. Everything should just use the function that takes the context. The other overloads should be removed, in a later PR.

…oss_compile_ctx

bdevorem

I've been using this branch for a few weeks actually and it has been working well, lgtm. I don't mind it being controlled by env vars, it has been easy and straightforward as a user, though I agree you could more cleanly control a lot more toggles if using compile options. There is a copilot comment I think is worth looking at, I had a comment about potential nonsense targets, and code coverage is red

bdevorem · 2026-05-28T16:10:59Z

 }

+template <class T>
+value to_value_target(const T& x)


is it possible to add test coverage for this?

bdevorem · 2026-05-28T16:14:01Z

+    {
+        if(target_name == "gpu" and not gpu_arch.empty())
+        {
+            return make_target(target_name,


what would happen if the values don't combine to be a valid target, simulated or otherwise? I'm not sure if that's even possible, but curious if there needs to be error checking here

It will throw down the line with a message like this:

Pass: gpu::compile_ops Unsupported architecture: gfx0

bdevorem · 2026-05-28T16:15:05Z

    return it->second;
 }

+target make_target(const std::string& name, const value& options)


test coverage here too

bdevorem

Licensing is red but I’ll approve now

pfultz2 · 2026-05-28T20:02:37Z


+MIGRAPHX_C_EXPORT migraphx_status migraphx_target_create_with_options(migraphx_target_t* target,
+                                                                      const char* name,
+                                                                      const char* options_json);


Just like migraphx_operation_create this should be variadic. So the user can do migraphx_target_create_with_options(&target, "gpu", "{gpu_arch: %s}", arch)

pfultz2 · 2026-05-28T20:04:57Z

 MIGRAPHX_EXPORT void register_target(const target& t);
 MIGRAPHX_EXPORT void unregister_target(const std::string& name);
 MIGRAPHX_EXPORT target make_target(const std::string& name);
+MIGRAPHX_EXPORT target make_target(const std::string& name, const value& options);


There should use std::initializer_list<std::pair<std::string, value>> and then add an template overload to handle the value type just like make_op.

pfultz2 · 2026-05-28T20:05:36Z

+        if(target_name == "gpu" and not gpu_arch.empty())
+        {
+            return make_target(target_name,
+                               value{{"gpu_arch", gpu_arch},


Remove value.

Suggested change

value{{"gpu_arch", gpu_arch},

{{"gpu_arch", gpu_arch},

pfultz2 · 2026-05-28T20:08:56Z

+{
+    if(options_json == nullptr or *options_json == '\0')
+        return make_target(name);
+    return make_target(name, from_json_string(options_json));


This needs to first call convert_to_json so users can pass in the keys without quotes: migraphx_target_create_with_options(&target, "gpu", "{num_cu: 32}"). This makes it consistent with migraphx_operation_create

pfultz2 · 2026-05-28T20:11:34Z

+inline namespace MIGRAPHX_INLINE_NS {
+namespace gpu {
+
+hipDeviceProp_t make_cross_compile_device_props(const std::string& arch_name, std::size_t cu_count)


Can we define our own device property struct instead of relying on hipDeviceProp_t? Hip could change it in the future.

…oss_compile_ctx

kahmed10 added 11 commits April 9, 2026 17:06

refactor device name functions

43b24e4

change to overload

aedf850

formatting

acf9c46

use context

06a966c

formatting

eed1afc

Merge branch 'develop' of https://github.com/ROCm/AMDMIGraphX into de…

6bcfc01

…vice_name_change

add context and compile changes

b95d1be

refactor

424f545

skip benchmarking if cross compile

0bbd889

formatting

9f53832

fix licensing

d10cd5a

fix cppcheck and tidy

cbde4eb

kahmed10 changed the title ~~Cross Compile Pt 2: add env variable to enable cross compilation~~ [AIMIGRAPHX-828] Cross Compile Pt 2: add env variable to enable cross compilation Apr 21, 2026

kahmed10 requested review from CharlieL7 and pfultz2 April 21, 2026 20:25

kahmed10 added 5 commits April 21, 2026 17:00

manual merge

976f9e0

fix merge

c288251

Merge branch 'develop' of https://github.com/ROCm/AMDMIGraphX into cr…

a56b279

…oss_compile_ctx

small refactor and fixing issues

5384040

formatting

e36c121

kahmed10 marked this pull request as ready for review April 23, 2026 20:09

kahmed10 requested a review from a team as a code owner April 23, 2026 20:09

Copilot AI review requested due to automatic review settings April 23, 2026 20:09

kahmed10 requested a review from causten as a code owner April 23, 2026 20:09

Copilot started reviewing on behalf of kahmed10 April 23, 2026 20:11 View session

Copilot AI reviewed Apr 23, 2026

View reviewed changes

pfultz2 reviewed Apr 24, 2026

View reviewed changes

kahmed10 added 10 commits May 13, 2026 14:22

Merge branch 'develop' of https://github.com/ROCm/AMDMIGraphX into cr…

afcf751

…oss_compile_ctx

formatting

2e94fd6

Merge branch 'develop' of https://github.com/ROCm/AMDMIGraphX into cr…

335b951

…oss_compile_ctx

fix tests

05087b9

fix license

2cd48e7

fix test

de57b4f

fix test

742bdaa

formatting

f906286

update test

5179e9e

Merge branch 'develop' of https://github.com/ROCm/AMDMIGraphX into cr…

0b9752b

…oss_compile_ctx

CharlieL7 approved these changes May 21, 2026

View reviewed changes

kahmed10 changed the title ~~[AIMIGRAPHX-828] Cross Compile Pt 2: add env variable to enable cross compilation~~ [AIMIGRAPHX-828] Cross Compile Pt 2: construct gpu target with gfx number to enable cross compilation May 21, 2026

CharlieL7 requested a review from bdevorem May 27, 2026 19:42

bdevorem requested changes May 28, 2026

View reviewed changes

kahmed10 added 3 commits May 28, 2026 13:04

manual merge

4891342

fix merge

7b98c5e

add more tests

6351529

kahmed10 requested a review from bdevorem May 28, 2026 18:45

bdevorem approved these changes May 28, 2026

View reviewed changes

kahmed10 added 2 commits May 28, 2026 13:56

Fix license

d883468

remove headers

e482533

pfultz2 requested changes May 28, 2026

View reviewed changes

kahmed10 added 5 commits June 1, 2026 13:18

update APIs

98d6088

formatting

8a04535

Merge branch 'develop' of https://github.com/ROCm/AMDMIGraphX into cr…

480fbab

…oss_compile_ctx

fix commented code

fa6df80

regenerate file

a9fc853

pfultz2 approved these changes Jun 2, 2026

View reviewed changes

causten merged commit 82079fc into develop Jun 2, 2026
36 of 38 checks passed

causten deleted the cross_compile_ctx branch June 2, 2026 18:02

		if(string_value_of(MIGRAPHX_GPU_ARCH{}).empty())
		this->finalize();

Conversation

kahmed10 commented Apr 17, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

Changelog Category

Uh oh!

codecov Bot commented Apr 17, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

kahmed10 commented Apr 23, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

bdevorem left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

bdevorem left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

kahmed10 commented Apr 17, 2026 •

edited

Loading

codecov Bot commented Apr 17, 2026 •

edited

Loading