[AIMIGRAPHX-828] Cross Compile Pt 2: add env variable 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 environment variable MIGRAPHX_GPU_ARCH which can be set to whatever is supported by the compiler.
when set, 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 60.00000% with 2 lines in your changes missing coverage. Please review.

Files with missing lines	Patch %	Lines
src/program.cpp	60.00%	2 Missing ⚠️

Additional details and impacted files

@@             Coverage Diff             @@
##           develop    #4795      +/-   ##
===========================================
- Coverage    92.49%   92.49%   -0.00%     
===========================================
  Files          583      583              
  Lines        29562    29565       +3     
===========================================
+ Hits         27343    27345       +2     
- Misses        2219     2220       +1

Files with missing lines	Coverage Δ
src/include/migraphx/program.hpp	`100.00% <ø> (ø)`
src/program.cpp	`72.00% <60.00%> (-0.02%)`	⬇️

🚀 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;


+      | 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.


+    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.

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

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[AIMIGRAPHX-828] Cross Compile Pt 2: add env variable to enable cross compilation#4795

[AIMIGRAPHX-828] Cross Compile Pt 2: add env variable to enable cross compilation#4795
kahmed10 wants to merge 17 commits intodevelopfrom
cross_compile_ctx

kahmed10 commented Apr 17, 2026

Uh oh!

codecov Bot commented Apr 17, 2026 •

edited

Loading

Uh oh!

kahmed10 commented Apr 23, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

pfultz2 Apr 24, 2026

Uh oh!

pfultz2 Apr 24, 2026

Uh oh!

pfultz2 Apr 24, 2026

Uh oh!

pfultz2 Apr 24, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

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

Conversation

kahmed10 commented Apr 17, 2026

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!

pfultz2 Apr 24, 2026

Choose a reason for hiding this comment

Uh oh!

pfultz2 Apr 24, 2026

Choose a reason for hiding this comment

Uh oh!

pfultz2 Apr 24, 2026

Choose a reason for hiding this comment

Uh oh!

pfultz2 Apr 24, 2026

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

codecov Bot commented Apr 17, 2026 •

edited

Loading