Problem Description
Summary
src/targets/gpu/lowering.cpp calls gpu::gfx_default_rocblas() unconditionally, but its declaration in src/targets/gpu/include/migraphx/gpu/device_name.hpp is guarded behind #if MIGRAPHX_USE_HIPBLASLT. Any build configured with -DMIGRAPHX_USE_HIPBLASLT=Off fails to compile.
This affects any GPU target with no hipBLASLt kernels at all (e.g. gfx900/gfx906 - Vega10/Vega20), where MIGRAPHX_USE_HIPBLASLT=Off is the only correct build configuration, not an edge case.
Error
/migraphx-src/src/targets/gpu/lowering.cpp:280:51: error: no member named 'gfx_default_rocblas' in namespace 'migraphx::gpu'
280 | not hipblaslt_supported() or gpu::gfx_default_rocblas()))
| ^~~~~~~~~~~~~~~~~~~
1 error generated.
Root cause
src/targets/gpu/include/migraphx/gpu/device_name.hpp:
#if MIGRAPHX_USE_HIPBLASLT
MIGRAPHX_GPU_EXPORT bool gfx_default_rocblas();
MIGRAPHX_GPU_EXPORT bool gfx_default_rocblas(const context& ctx);
#endif
MIGRAPHX_GPU_EXPORT bool hipblaslt_supported();
MIGRAPHX_GPU_EXPORT bool hipblaslt_supported(const context& ctx);
src/targets/gpu/lowering.cpp (no matching guard around the call):
if(not has_fp8_inputs and
((string_value_of(MIGRAPHX_SET_GEMM_PROVIDER{}) == "rocblas") or
not hipblaslt_supported() or gpu::gfx_default_rocblas()))
hipblaslt_supported() is itself declared/defined unconditionally and correctly returns a hardcoded false when the flag is off (device_name.cpp):
bool hipblaslt_supported()
{
#if !MIGRAPHX_USE_HIPBLASLT
return false;
#else
return hipblaslt_supported_impl(get_gfx_name(get_device_name()));
#endif
}
So the asymmetry is specifically in gfx_default_rocblas()'s declaration being conditionally compiled out while its only real caller (lowering.cpp) is not.
Reproduction
git clone --depth 1 https://github.com/ROCm/AMDMIGraphX.git
cd AMDMIGraphX
cmake -B build -G Ninja -DGPU_TARGETS=gfx900 -DMIGRAPHX_USE_HIPBLASLT=Off -DMIGRAPHX_USE_COMPOSABLEKERNEL=Off -DCMAKE_INSTALL_PREFIX=/opt/rocm
cmake --build build --target install
Fails identically on develop (current HEAD as of this report) and on release/rocm-rel-7.14.
You can see my failed CI run here: https://github.com/Schaka/rocm-migraphx-ort-builder/actions/runs/30455557112/job/90588160497
Suggested fix
Since not hipblaslt_supported() already evaluates unconditionally to true whenever MIGRAPHX_USE_HIPBLASLT is off, the whole or chain is already true regardless of gfx_default_rocblas()'s value in that configuration - dropping the redundant call when the flag is off is a semantics-preserving fix, not a behavior change:
- not hipblaslt_supported() or gpu::gfx_default_rocblas()))
+ not hipblaslt_supported()
+ #if MIGRAPHX_USE_HIPBLASLT
+ or gpu::gfx_default_rocblas()
+ #endif
+ ))
(or equivalently, wrap the declaration usage with the same #if MIGRAPHX_USE_HIPBLASLT guard used in the header, providing a false fallback for the dropped disjunct - either form compiles and preserves current behavior.)
Environment
- Base:
rocm/dev-ubuntu-26.04:7.14.0-full
- Target:
-DGPU_TARGETS=gfx900 (also reproduces on gfx906)
- Compiler:
/opt/rocm/llvm/bin/clang++ (ROCm 7.14.0 bundled clang)
Operating System
Ubuntu 26.04 (Docker)
CPU
Ultra 7 270k locally and whatever GitHub uses in CI
GPU
Other
Other
Irrelevant, happens at compile time
ROCm Version
ROCm 7.14
Steps to Reproduce
No response
(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support
No response
Additional Information
No response
Problem Description
Summary
src/targets/gpu/lowering.cppcallsgpu::gfx_default_rocblas()unconditionally, but its declaration insrc/targets/gpu/include/migraphx/gpu/device_name.hppis guarded behind#if MIGRAPHX_USE_HIPBLASLT. Any build configured with-DMIGRAPHX_USE_HIPBLASLT=Offfails to compile.This affects any GPU target with no hipBLASLt kernels at all (e.g. gfx900/gfx906 - Vega10/Vega20), where
MIGRAPHX_USE_HIPBLASLT=Offis the only correct build configuration, not an edge case.Error
Root cause
src/targets/gpu/include/migraphx/gpu/device_name.hpp:src/targets/gpu/lowering.cpp(no matching guard around the call):hipblaslt_supported()is itself declared/defined unconditionally and correctly returns a hardcodedfalsewhen the flag is off (device_name.cpp):So the asymmetry is specifically in
gfx_default_rocblas()'s declaration being conditionally compiled out while its only real caller (lowering.cpp) is not.Reproduction
Fails identically on
develop(current HEAD as of this report) and onrelease/rocm-rel-7.14.You can see my failed CI run here: https://github.com/Schaka/rocm-migraphx-ort-builder/actions/runs/30455557112/job/90588160497
Suggested fix
Since
not hipblaslt_supported()already evaluates unconditionally totruewheneverMIGRAPHX_USE_HIPBLASLTis off, the wholeorchain is alreadytrueregardless ofgfx_default_rocblas()'s value in that configuration - dropping the redundant call when the flag is off is a semantics-preserving fix, not a behavior change:(or equivalently, wrap the declaration usage with the same
#if MIGRAPHX_USE_HIPBLASLTguard used in the header, providing afalsefallback for the dropped disjunct - either form compiles and preserves current behavior.)Environment
rocm/dev-ubuntu-26.04:7.14.0-full-DGPU_TARGETS=gfx900(also reproduces ongfx906)/opt/rocm/llvm/bin/clang++(ROCm 7.14.0 bundled clang)Operating System
Ubuntu 26.04 (Docker)
CPU
Ultra 7 270k locally and whatever GitHub uses in CI
GPU
Other
Other
Irrelevant, happens at compile time
ROCm Version
ROCm 7.14
Steps to Reproduce
No response
(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support
No response
Additional Information
No response