Skip to content

Check if workaround deduction guide is needed#4763

Open
pfultz2 wants to merge 3 commits intodevelopfrom
hip-workaround-deduction-guide
Open

Check if workaround deduction guide is needed#4763
pfultz2 wants to merge 3 commits intodevelopfrom
hip-workaround-deduction-guide

Conversation

@pfultz2
Copy link
Copy Markdown
Collaborator

@pfultz2 pfultz2 commented Apr 9, 2026

Motivation

The deduction guides for aggregates constructors have been broken on hip as it doesnt propagate the device target from the constructor. This has been fixed on a newer compiler and the workaround is deprecated behaviour. For now, detect when this behaviour works and adjust the workaround. The MIGRAPHX_AUTO_DEDUCE can be removed on a newer version of the compiler.

Technical Details

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.

Copilot AI review requested due to automatic review settings April 9, 2026 13:15
@pfultz2 pfultz2 requested a review from causten as a code owner April 9, 2026 13:15
Copy link
Copy Markdown
Contributor

Copilot AI left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull request overview

Updates MIGraphX’s HIP kernel compilation path to conditionally enable a workaround for broken class template argument deduction (CTAD) deduction guides on certain HIP/Clang toolchains.

Changes:

  • Make MIGRAPHX_AUTO_DEDUCE conditional on a new MIGRAPHX_WORKAROUND_BROKEN_DEDUCTION_GUIDE macro.
  • Introduce hip_can_compile(src, flags) to support compiling arbitrary probe sources (and refactor hip_has_flags to use it).
  • Add a compile-time probe in HIP code object compilation to decide whether to define MIGRAPHX_WORKAROUND_BROKEN_DEDUCTION_GUIDE, and extend warning filtering.

Reviewed changes

Copilot reviewed 4 out of 4 changed files in this pull request and generated 2 comments.

File Description
src/targets/gpu/kernels/include/migraphx/kernels/types.hpp Makes deduction guide emission conditional (workaround vs non-workaround form).
src/targets/gpu/include/migraphx/gpu/compile_hip.hpp Exposes new hip_can_compile helper in the GPU compile API.
src/targets/gpu/compile_hip.cpp Implements hip_can_compile and refactors hip_has_flags to call it.
src/targets/gpu/compile_hip_code_object.cpp Adds a compile probe to decide if the deduction-guide workaround macro should be defined; adds a warning suppression probe.
Comments suppressed due to low confidence (1)

src/targets/gpu/compile_hip.cpp:381

  • hip_can_compile hardcodes arch = "gfx900", so the result depends on whether the toolchain supports that specific target rather than whether the provided source/flags are generally compilable. Since this is now a public/utility API (and is used to gate behavior), consider accepting an arch argument (or using the current device arch / a compiler-supported default like native) to avoid false failures caused by an unsupported probe arch.
bool hip_can_compile(const std::string& src, const std::vector<std::string>& flags)
{
    src_file input{"main.cpp", src};
    std::vector<src_file> srcs = {input};

    try
    {
        std::string arch = "gfx900";
        compile_hip_src(srcs, flags, arch, true);
        return true;

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

template<class T>
S(T) -> S<T>;

__attribute__((device)) auto f()
Copy link

Copilot AI Apr 9, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The probe uses __attribute__((device)), which may be ignored/treated as an unknown attribute by Clang in HIP mode (often only emitting a warning). If that happens, this test can compile even on compilers where the deduction guide is still broken, causing the workaround to be disabled incorrectly. Prefer using HIP’s __device__ (or a __global__ kernel) so the code is unquestionably compiled in device context.

Suggested change
__attribute__((device)) auto f()
__device__ auto f()

Copilot uses AI. Check for mistakes.

)__migraphx__";

static bool can_compile = hip_can_compile(test, {"-std=c++17"});
Copy link

Copilot AI Apr 9, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

hip_can_compile(test, {"-std=c++17"}) allows the probe to succeed even if the compiler only emits warnings (e.g., for unknown attributes/options), which can lead to a false “works” result and disable the workaround. Consider passing -Werror (and any other minimal flags needed) to make this feature-detection conservative (false positives are OK; false negatives are risky).

Suggested change
static bool can_compile = hip_can_compile(test, {"-std=c++17"});
static bool can_compile = hip_can_compile(test, {"-std=c++17", "-Werror"});

Copilot uses AI. Check for mistakes.
@codecov
Copy link
Copy Markdown

codecov bot commented Apr 9, 2026

Codecov Report

✅ All modified and coverable lines are covered by tests.

Additional details and impacted files
@@             Coverage Diff             @@
##           develop    #4763      +/-   ##
===========================================
- Coverage    92.37%   92.32%   -0.04%     
===========================================
  Files          581      583       +2     
  Lines        28703    29334     +631     
===========================================
+ Hits         26512    27082     +570     
- Misses        2191     2252      +61     

see 34 files with indirect coverage changes

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

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.

3 participants