Check if workaround deduction guide is needed#4763
Conversation
There was a problem hiding this comment.
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_DEDUCEconditional on a newMIGRAPHX_WORKAROUND_BROKEN_DEDUCTION_GUIDEmacro. - Introduce
hip_can_compile(src, flags)to support compiling arbitrary probe sources (and refactorhip_has_flagsto 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_compilehardcodesarch = "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 anarchargument (or using the current device arch / a compiler-supported default likenative) 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() |
There was a problem hiding this comment.
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.
| __attribute__((device)) auto f() | |
| __device__ auto f() |
|
|
||
| )__migraphx__"; | ||
|
|
||
| static bool can_compile = hip_can_compile(test, {"-std=c++17"}); |
There was a problem hiding this comment.
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).
| static bool can_compile = hip_can_compile(test, {"-std=c++17"}); | |
| static bool can_compile = hip_can_compile(test, {"-std=c++17", "-Werror"}); |
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 🚀 New features to boost your workflow:
|
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_DEDUCEcan be removed on a newer version of the compiler.Technical Details
Changelog Category
Add a
CHANGELOG.mdentry for any option other thanNot Applicable