Error check to avoid null config dereference #4319

harakas · 2025-09-25T10:47:06Z

Motivation

Give a more proper error instead of abort on std::bad_alloc.

Technical Details

I'm trying to build migraphx on top of TheRock. It does not work and aborts with std::bad_alloc. I traced this to a null std::optional dereference (after recompiling with -D_GLIBCXX_ASSERTIONS):

# /opt/rocm/bin/migraphx-driver compile --gpu yolov8n_320.onnx                   
Running [ MIGraphX Version: 2.14.0.20250912-17-96-g03f38cf75-dirty ]: /opt/rocm/bin/migraphx-driver compile --gpu yolov8n_320.onnx
[2025-09-25 10:24:25]
Reading: yolov8n_320.onnx
Compiling ... 
/usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/optional:484: const _Tp &std::_Optional_base_impl<migraphx::gpu::tuning_config, std::_Optional_base<migraphx::gpu::tuning_config>>::_M_get() const [_Tp = migraphx::gpu::tuning_config, _Dp = std::_Optional_base<migraphx::gpu::tuning_config>]: Assertion 'this->_M_is_engaged()' failed.
Aborted (core dumped)

The fix is to check for null config. This now results in a migraphx error:

# /opt/rocm/bin/migraphx-driver compile --gpu yolov8n_320.onnx
Running [ MIGraphX Version: 2.14.0.20250912-17-96-g03f38cf75-dirty ]: /opt/rocm/bin/migraphx-driver compile --gpu yolov8n_320.onnx
[2025-09-25 10:25:41]
Reading: yolov8n_320.onnx
Compiling ... 
terminate called after throwing an instance of 'migraphx::version_2_14_0::exception'
  what():  /AMDMIGraphX/src/targets/gpu/compile_ops.cpp:212: benchmark: Kernel without config for concat
Aborted (core dumped)

I have no idea what the error means though. Migraphx works just fine when I compile the same way against a stock ROCm install. So I suppose TheRock must be doing something strange.

Changelog Category

- 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

Pull Request Overview

This PR adds error checking to prevent a null dereference when accessing a kernel configuration, replacing a confusing std::bad_alloc abort with a clearer MIGraphX exception. The change provides better error diagnostics when no configuration is found for a kernel operation.

Adds null check for config before dereferencing
Throws descriptive MIGraphX exception instead of allowing null dereference abort

_{Tip: Customize your code reviews with copilot-instructions.md. Create the file or learn how to get started.}

pfultz2 · 2025-09-28T21:11:17Z

src/targets/gpu/compile_ops.cpp

        if(results.size() == 1)
        {
+            if(not config)
+                MIGRAPHX_THROW("Kernel without config for " + preop.name());


A config is not required when there is only one kernel.

pfultz2 · 2025-09-28T21:16:07Z

Do you have a backtrace for the debug build as this does not seem to be the right place to check for the null config?

harakas · 2025-09-29T15:13:27Z

Here's a backtrace @pfultz2

...
[Thread 0x7cfac8df7700 (LWP 12887) exited]
[New Thread 0x7cfac8df7700 (LWP 12888)]
[Thread 0x7cfac8df7700 (LWP 12888) exited]
[Thread 0x7cfa3a8a8700 (LWP 11844) exited]
[New Thread 0x7cfa3a0a7700 (LWP 12889)]
/usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/optional:484: const _Tp &std::_Optional_base_impl<migraphx::gpu::tuning_config, std::_Optional_base<migraphx::gpu::tuning_config>>::_M_get() const [_Tp = migraphx::gpu::tuning_config, _Dp = std::_Optional_base<migraphx::gpu::tuning_config>]: Assertion 'this->_M_is_engaged()' failed.
--Type <RET> for more, q to quit, c to continue without paging--

Thread 1 "migraphx-driver" received signal SIGABRT, Aborted.
__GI_raise (sig=sig@entry=6) at ../sysdeps/unix/sysv/linux/raise.c:51
51      ../sysdeps/unix/sysv/linux/raise.c: No such file or directory.
(gdb) bt
#0  __GI_raise (sig=sig@entry=6) at ../sysdeps/unix/sysv/linux/raise.c:51
#1  0x00007cfadad9e7f1 in __GI_abort () at abort.c:79
#2  0x00007cfadb7ec649 in std::__glibcxx_assert_fail(char const*, int, char const*, char const*) () from /usr/lib/x86_64-linux-gnu/libstdc++.so.6
#3  0x00007cfb088c7ba4 in std::_Optional_base_impl<migraphx::version_2_14_0::gpu::tuning_config, std::_Optional_base<migraphx::version_2_14_0::gpu::tuning_config, false, false> >::_M_get (this=0x5e45b034b0c8) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/optional:484
#4  0x00007cfb088c7b55 in std::optional<migraphx::version_2_14_0::gpu::tuning_config>::operator-> (this=0x5e45b034b0c8)
    at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/optional:962
#5  0x00007cfb088c38c9 in migraphx::version_2_14_0::gpu::compile_plan::print_modules[abi:cxx11]() const (this=0x5e45b034b0a0)
    at /workspace/AMDMIGraphX/src/targets/gpu/compile_ops.cpp:186
#6  0x00007cfb088c9505 in migraphx::version_2_14_0::gpu::compile_plan::benchmark (this=0x5e45b034b0a0) at /workspace/AMDMIGraphX/src/targets/gpu/compile_ops.cpp:214
#7  0x00007cfb088c3009 in migraphx::version_2_14_0::gpu::compile_plan::replace (this=0x5e45b034b0a0, m=...) at /workspace/AMDMIGraphX/src/targets/gpu/compile_ops.cpp:289
#8  0x00007cfb088c04b5 in migraphx::version_2_14_0::gpu::compile_manager::compile (this=0x7ffce165d0a0, m=...)
    at /workspace/AMDMIGraphX/src/targets/gpu/compile_ops.cpp:335
#9  0x00007cfb088bfea1 in migraphx::version_2_14_0::gpu::compile_ops::apply (this=0x5e45b04894f8, m=...) at /workspace/AMDMIGraphX/src/targets/gpu/compile_ops.cpp:359
#10 0x00007cfb08c07c6d in migraphx::version_2_14_0::detail::module_pass_manager_apply<migraphx::version_2_14_0::gpu::compile_ops> (x=..., mpm=...)
    at /workspace/AMDMIGraphX/src/include/migraphx/pass.hpp:69
#11 0x00007cfb08c07c2d in migraphx::version_2_14_0::detail::module_pass_manager_apply<migraphx::version_2_14_0::gpu::compile_ops> (x=..., mpm=...)
    at /workspace/AMDMIGraphX/src/include/migraphx/pass.hpp:80
#12 0x00007cfb08c07c02 in migraphx::version_2_14_0::pass::private_detail_te_default_apply<migraphx::version_2_14_0::gpu::compile_ops const&> (private_detail_te_self=...,
    mpm=...) at /workspace/AMDMIGraphX/src/include/migraphx/pass.hpp:115
#13 0x00007cfb08c07874 in migraphx::version_2_14_0::pass::private_detail_te_handle_type<migraphx::version_2_14_0::gpu::compile_ops>::apply (this=0x5e45b04894f0, mpm=...)
    at /workspace/AMDMIGraphX/src/include/migraphx/pass.hpp:296
#14 0x00007cfae173aeaa in migraphx::version_2_14_0::pass::apply (this=0x5e45b0489d00, mpm=...) at /workspace/AMDMIGraphX/src/include/migraphx/pass.hpp:236
#15 0x00007cfae173a0c0 in migraphx::version_2_14_0::module_pm::run_pass (this=0x7ffce165d380, p=...) at /workspace/AMDMIGraphX/src/pass_manager.cpp:149
#16 0x00007cfae1738f3c in migraphx::version_2_14_0::run_passes (prog=..., root_mod=0x5e45b03081b8, passes=std::vector of length 87, capacity 87 = {...}, trace=...)
    at /workspace/AMDMIGraphX/src/pass_manager.cpp:187
#17 0x00007cfae17391b1 in migraphx::version_2_14_0::run_passes (prog=..., passes=std::vector of length 87, capacity 87 = {...}, trace=...)
    at /workspace/AMDMIGraphX/src/pass_manager.cpp:205
#18 0x00007cfae174c3fd in migraphx::version_2_14_0::program::compile (this=0x7ffce165dfe0, t=..., options=...) at /workspace/AMDMIGraphX/src/program.cpp:326
#19 0x00005e458a919e12 in migraphx::driver::version_2_14_0::compiler::compile (this=0x7ffce165e208) at /workspace/AMDMIGraphX/src/driver/main.cpp:666
#20 0x00005e458a918e89 in migraphx::driver::version_2_14_0::compile::run (this=0x7ffce165e208) at /workspace/AMDMIGraphX/src/driver/main.cpp:777
#21 0x00005e458a918dc8 in migraphx::driver::version_2_14_0::run_command<migraphx::driver::version_2_14_0::compile> (
    exe_name="/workspace/AMDMIGraphX/build/bin/migraphx-driver", args=std::vector of length 2, capacity 2 = {...}, add_help=true)
    at /workspace/AMDMIGraphX/src/driver/command.hpp:83
#22 0x00005e458a9187e6 in migraphx::driver::version_2_14_0::auto_register_command<migraphx::driver::version_2_14_0::compile>()::{lambda(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > > const&)#1}::operator()(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > > const&) const (this=0x5e45b0302498, exe_name="/workspace/AMDMIGraphX/build/bin/migraphx-driver",
    args=std::vector of length 2, capacity 2 = {...}) at /workspace/AMDMIGraphX/src/driver/command.hpp:91
#23 0x00005e458a9187b5 in std::__invoke_impl<void, migraphx::driver::version_2_14_0::auto_register_command<migraphx::driver::version_2_14_0::compile>()::{lambda(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > > const&)#1}&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::--Type <RET> for more, q to quit, c to continue without paging--
char_traits<char>, std::allocator<char> > > > >(std::__invoke_other, migraphx::driver::version_2_14_0::auto_register_command<migraphx::driver::version_2_14_0::compile>()::{lambda(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > > const&)#1}&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > >&&) (__f=..., __args=..., __args=...)
    at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/invoke.h:61
#24 0x00005e458a918755 in std::__invoke_r<void, migraphx::driver::version_2_14_0::auto_register_command<migraphx::driver::version_2_14_0::compile>()::{lambda(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > > const&)#1}&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > > >(migraphx::driver::version_2_14_0::auto_register_command<migraphx::driver::version_2_14_0::compile>()::{lambda(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > > const&)#1}&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > >&&) (__fn=..., __args=..., __args=...) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/invoke.h:111
#25 0x00005e458a91867d in std::_Function_handler<void (std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > >), migraphx::driver::version_2_14_0::auto_register_command<migraphx::driver::version_2_14_0::compile>()::{lambda(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > > const&)#1}>::_M_invoke(std::_Any_data const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > >&&) (__functor=..., __args=..., __args=...) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/std_function.h:290
#26 0x00005e458a8ca8b2 in std::function<void (std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > >)>::operator()(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > >) const (this=0x5e45b0302498,
    __args=std::vector of length 2, capacity 2 = {...}, __args=std::vector of length 2, capacity 2 = {...})
    at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/std_function.h:591
#27 0x00005e458a8c8f81 in main (argc=4, argv=0x7ffce165eb78, envp=0x5e45aff74a80) at /workspace/AMDMIGraphX/src/driver/main.cpp:1024
(gdb)

Edit:

(gdb) frame 6
#6  0x00007f2bc9e4b505 in migraphx::version_2_14_0::gpu::compile_plan::benchmark (this=0x5df7373b6b80) at /workspace/AMDMIGraphX/src/targets/gpu/compile_ops.cpp:214
214                     MIGRAPHX_THROW("No valid tuned compilation for " + preop.name() + " with " +
(gdb)

harakas · 2025-09-29T19:46:59Z

There seem to be more std::bad_alloc exceptions that are blanket caught by compile_plan::insert_compiles (but are not caught by -D_GLIBCXX_ASSERTIONS):

# MIGRAPHX_TRACE_BENCHMARKING=1 /workspace/AMDMIGraphX/build/bin/migraphx-driver compile --gpu yolov8n_320.onnx
Running [ MIGraphX Version: 2.14.0.20250912-17-102-g18f81a6c6-dirty ]: /workspace/AMDMIGraphX/build/bin/migraphx-driver compile --gpu yolov8n_320.onnx
[2025-09-29 15:42:36]
Reading: yolov8n_320.onnx
Compiling ... 
Exception in contiguous: std::bad_alloc
Exception in concat: std::bad_alloc
Exception in concat: std::bad_alloc
Exception in concat: std::bad_alloc
Exception in concat: std::bad_alloc
Exception in concat: std::bad_alloc
Exception in concat: std::bad_alloc
Exception in concat: std::bad_alloc
Exception in concat: std::bad_alloc
Exception in concat: std::bad_alloc
Exception in fused_reduce: std::bad_alloc
Exception in layout: std::bad_alloc
Exception in concat: std::bad_alloc
Exception in contiguous: std::bad_alloc
Exception in concat: std::bad_alloc
Exception in pointwise: std::bad_alloc
Exception in pointwise: std::bad_alloc
Exception in concat: std::bad_alloc
Exception in fused_concat: std::bad_alloc
Exception in concat: std::bad_alloc
Exception in concat: std::bad_alloc
Exception in concat: std::bad_alloc
Exception in concat: std::bad_alloc
Exception in concat: std::bad_alloc
Exception in concat: std::bad_alloc
Benchmarking gpu::mlir_op: 25 configs
Fastest solution: v3:64,64,128,4,2,2,1,1,2
Benchmarking gpu::mlir_op: 25 configs
Fastest solution: v3:64,64,128,4,2,2,1,1,2
Benchmarking gpu::mlir_op: 25 configs
Fastest solution: v3:64,32,64,8,4,2,1,1,2
Benchmarking gpu::mlir_op: 25 configs
Fastest solution: v3:64,64,128,4,2,2,1,1,2
Benchmarking concat: 1 configs
/usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/optional:484: const _Tp &std::_Optional_base_impl<migraphx::gpu::tuning_config, std::_Optional_base<migraphx::gpu::tuning_config>>::_M_get() const [_Tp = migraphx::gpu::tuning_config, _Dp = std::_Optional_base<migraphx::gpu::tuning_config>]: Assertion 'this->_M_is_engaged()' failed.

Commenting out the blanket try/catches yields:

#0  __GI_raise (sig=sig@entry=6) at ../sysdeps/unix/sysv/linux/raise.c:51
#1  0x000074217fd5c7f1 in __GI_abort () at abort.c:79
#2  0x000074218076e646 in ?? () from /usr/lib/x86_64-linux-gnu/libstdc++.so.6
#3  0x000074218077ff06 in ?? () from /usr/lib/x86_64-linux-gnu/libstdc++.so.6
#4  0x000074218077ff71 in std::terminate() () from /usr/lib/x86_64-linux-gnu/libstdc++.so.6
#5  0x0000742180780207 in __cxa_rethrow () from /usr/lib/x86_64-linux-gnu/libstdc++.so.6
#6  0x00007421ad8cb350 in std::__do_uninit_copy<__gnu_cxx::__normal_iterator<migraphx::version_2_14_0::src_file const*, std::vector<migraphx::version_2_14_0::src_file, std::allocator<migraphx::version_2_14_0::src_file> > >, migraphx::version_2_14_0::gpu::hiprtc_src_file*> (__first=..., __last=..., __result=0x7420b8612f50) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/stl_uninitialized.h:126
#7  0x00007421ad8cb2b5 in std::__uninitialized_copy<false>::__uninit_copy<__gnu_cxx::__normal_iterator<migraphx::version_2_14_0::src_file const*, std::vector<migraphx::version_2_14_0::src_file, std::allocator<migraphx::version_2_14_0::src_file> > >, migraphx::version_2_14_0::gpu::hiprtc_src_file*> (__first=..., __last=..., __result=0x7420b8612f50) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/stl_uninitialized.h:137
#8  0x00007421ad8cb26d in std::uninitialized_copy<__gnu_cxx::__normal_iterator<migraphx::version_2_14_0::src_file const*, std::vector<migraphx::version_2_14_0::src_file, std::allocator<migraphx::version_2_14_0::src_file> > >, migraphx::version_2_14_0::gpu::hiprtc_src_file*> (__first=..., __last=..., __result=0x7420b8612f50) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/stl_uninitialized.h:184
#9  0x00007421ad8cb109 in std::__uninitialized_copy_a<__gnu_cxx::__normal_iterator<migraphx::version_2_14_0::src_file const*, std::vector<migraphx::version_2_14_0::src_file, std::allocator<migraphx::version_2_14_0::src_file> > >, migraphx::version_2_14_0::gpu::hiprtc_src_file*, migraphx::version_2_14_0::gpu::hiprtc_src_file> (__first=..., __last=..., __result=0x7420b8612f50) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/stl_uninitialized.h:373
#10 0x00007421ad8caf62 in std::vector<migraphx::version_2_14_0::gpu::hiprtc_src_file, std::allocator<migraphx::version_2_14_0::gpu::hiprtc_src_file> >::_M_range_initialize<__gnu_cxx::__normal_iterator<migraphx::version_2_14_0::src_file const*, std::vector<migraphx::version_2_14_0::src_file, std::allocator<migraphx::version_2_14_0::src_file> > > > (this=0x7420e1a9c1d8, __first=..., __last=...) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/stl_vector.h:1692
#11 0x00007421ad8c50c2 in std::vector<migraphx::version_2_14_0::gpu::hiprtc_src_file, std::allocator<migraphx::version_2_14_0::gpu::hiprtc_src_file> >::vector<__gnu_cxx::__normal_iterator<migraphx::version_2_14_0::src_file const*, std::vector<migraphx::version_2_14_0::src_file, std::allocator<migraphx::version_2_14_0::src_file> > >, void> (this=0x7420e1a9c1d8, __first=..., __last=..., __a=...) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/stl_vector.h:708
#12 0x00007421ad8c0ddc in migraphx::version_2_14_0::gpu::compile_hip_src (srcs=..., params=..., arch=..., quiet=false) at /workspace/AMDMIGraphX/src/targets/gpu/compile_hip.cpp:234
#13 0x00007421ad8cd6df in migraphx::version_2_14_0::gpu::compile_hip_code_object (ctx=..., content=..., options=...) at /workspace/AMDMIGraphX/src/targets/gpu/compile_hip_code_object.cpp:205
#14 0x00007421adc07842 in migraphx::version_2_14_0::gpu::pointwise_compiler::compile_op (this=0x58f1f6b9e8c8, ctx=..., inputs=..., v=...) at /workspace/AMDMIGraphX/src/targets/gpu/jit/pointwise.cpp:100
#15 0x00007421adc067ed in migraphx::version_2_14_0::gpu::pointwise_compiler::compile (this=0x58f1f6b9e8c8, ctx=..., ins=..., op=...) at /workspace/AMDMIGraphX/src/targets/gpu/jit/pointwise.cpp:107
#16 0x00007421adc06560 in migraphx::version_2_14_0::gpu::compiler<migraphx::version_2_14_0::gpu::pointwise_compiler>::invoke_compile<migraphx::version_2_14_0::gpu::pointwise_compiler> (this=0x58f1f6b9e8c8, ctx=..., ins=..., op=..., solution=...) at /workspace/AMDMIGraphX/src/targets/gpu/include/migraphx/gpu/compiler.hpp:203
#17 0x00007421adc0648a in _ZZN8migraphx14version_2_14_03gpu17register_compilerINS1_18pointwise_compilerEEEvvENKUlDpOT_E_clIJRNS1_7contextENS0_15instruction_refENS0_9operationERKNS0_5valueEEEEDaS6_ (this=0x58f1f6b9e8c8, xs=..., xs=..., xs=..., xs=...) at /workspace/AMDMIGraphX/src/targets/gpu/include/migraphx/gpu/compiler.hpp:151
#18 0x00007421adc06400 in _ZSt13__invoke_implIN8migraphx14version_2_14_03gpu16compiler_replaceERZNS2_17register_compilerINS2_18pointwise_compilerEEEvvEUlDpOT_E_JRNS2_7contextENS1_15instruction_refENS1_9operationERKNS1_5valueEEET_St14__invoke_otherOT0_DpOT1_ (__f=..., __args=..., __args=..., __args=..., __args=...) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/invoke.h:61
#19 0x00007421adc06380 in _ZSt10__invoke_rIN8migraphx14version_2_14_03gpu16compiler_replaceERZNS2_17register_compilerINS2_18pointwise_compilerEEEvvEUlDpOT_E_JRNS2_7contextENS1_15instruction_refENS1_9operationERKNS1_5valueEEENSt9enable_ifIX16is_invocable_r_vIT_T0_DpT1_EESJ_E4typeEOSK_DpOSL_ (__fn=..., __args=..., __args=..., __args=..., __args=...) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/invoke.h:114
#20 0x00007421adc06280 in _ZNSt17_Function_handlerIFN8migraphx14version_2_14_03gpu16compiler_replaceERNS2_7contextENS1_15instruction_refENS1_9operationERKNS1_5valueEEZNS2_17register_compilerINS2_18pointwise_compilerEEEvvEUlDpOT_E_E9_M_invokeERKSt9_Any_dataS5_OS6_OS7_SA_ (__functor=..., __args=..., __args=..., __args=..., __args=...) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/std_function.h:290
#21 0x00007421ad8e0875 in std::function<migraphx::version_2_14_0::gpu::compiler_replace (migraphx::version_2_14_0::gpu::context&, migraphx::version_2_14_0::instruction_ref, migraphx::version_2_14_0::operation, migraphx::version_2_14_0::value const&)>::operator()(migraphx::version_2_14_0::gpu::context&, migraphx::version_2_14_0::instruction_ref, migraphx::version_2_14_0::operation, migraphx::version_2_14_0::value const&) const (this=0x58f1f6b9e8c8, __args=..., __args=..., __args=..., __args=...) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/std_function.h:591
#22 0x00007421ad8de9fa in migraphx::version_2_14_0::gpu::compile (ctx=..., ins=..., op=..., solution=...) at /workspace/AMDMIGraphX/src/targets/gpu/compiler.cpp:60
#23 0x00007421ad88433d in migraphx::version_2_14_0::gpu::compile_plan::insert_compiles<std::vector<std::function<void ()>, std::allocator<std::function<void ()> > > >(std::vector<std::function<void ()>, std::allocator<std::function<void ()> > >&, migraphx::version_2_14_0::value const&, unsigned long)::{lambda()#1}::operator()() const (this=0x58f1f6eae410) at /workspace/AMDMIGraphX/src/targets/gpu/compile_ops.cpp:113
#24 0x00007421ad8842e5 in std::__invoke_impl<void, migraphx::version_2_14_0::gpu::compile_plan::insert_compiles<std::vector<std::function<void ()>, std::allocator<std::function<void ()> > > >(std::vector<std::function<void ()>, std::allocator<std::function<void ()> > >&, migraphx::version_2_14_0::value const&, unsigned long)::{lambda()#1}&>(std::__invoke_other, migraphx::version_2_14_0::gpu::compile_plan::insert_compiles<std::vector<std::function<void ()>, std::allocator<std::function<void ()> > > >(std::vector<std::function<void ()>, std::allocator<std::function<void ()> > >&, migraphx::version_2_14_0::value const&, unsigned long)::{lambda()#1}&) (__f=...) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/invoke.h:61
#25 0x00007421ad8842a5 in std::__invoke_r<void, migraphx::version_2_14_0::gpu::compile_plan::insert_compiles<std::vector<std::function<void ()>, std::allocator<std::function<void ()> > > >(std::vector<std::function<void ()>, std::allocator<std::function<void ()> > >&, migraphx::version_2_14_0::value const&, unsigned long)::{lambda()#1}&>(migraphx::version_2_14_0::gpu::compile_plan::insert_compiles<std::vector<std::function<void ()>, std::allocator<std::function<void ()> > > >(std::vector<std::function<void ()>, std::allocator<std::function<void ()> > >&, migraphx::version_2_14_0::value const&, unsigned long)::{lambda()#1}&) (__fn=...) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/invoke.h:111
#26 0x00007421ad88411d in std::_Function_handler<void (), migraphx::version_2_14_0::gpu::compile_plan::insert_compiles<std::vector<std::function<void ()>, std::allocator<std::function<void ()> > > >(std::vector<std::function<void ()>, std::allocator<std::function<void ()> > >&, migraphx::version_2_14_0::value const&, unsigned long)::{lambda()#1}>::_M_invoke(std::_Any_data const&) (__functor=...) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/std_function.h:290
#27 0x00007421ad88656e in std::function<void ()>::operator()() const (this=0x58f1f6ee7360) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/std_function.h:591
#28 0x00007421ad8864b8 in migraphx::version_2_14_0::gpu::compile_manager::compile(migraphx::version_2_14_0::module&)::{lambda(auto:1)#1}::operator()<unsigned long>(unsigned long) const (this=0x7420e1a9da18, i=533) at /workspace/AMDMIGraphX/src/targets/gpu/compile_ops.cpp:328
#29 0x00007421ad886411 in migraphx::version_2_14_0::thread_invoke<migraphx::version_2_14_0::gpu::compile_manager::compile(migraphx::version_2_14_0::module&)::{lambda(auto:1)#1}>(unsigned long, unsigned long, migraphx::version_2_14_0::gpu::compile_manager::compile(migraphx::version_2_14_0::module&)::{lambda(auto:1)#1}) (i=533, f=...) at /workspace/AMDMIGraphX/src/include/migraphx/simple_par_for.hpp:62
#30 0x00007421ad8869bc in migraphx::version_2_14_0::simple_par_for_impl<migraphx::version_2_14_0::gpu::compile_manager::compile(migraphx::version_2_14_0::module&)::{lambda(auto:1)#1}>(unsigned long, unsigned long, migraphx::version_2_14_0::gpu::compile_manager::compile(migraphx::version_2_14_0::module&)::{lambda(auto:1)#1})::{lambda()#1}::operator()() const::{lambda()#1}::operator()() const (this=0x58f1f6eb8508) at /workspace/AMDMIGraphX/src/include/migraphx/simple_par_for.hpp:90
#31 0x00007421ad886945 in std::__invoke_impl<void, migraphx::version_2_14_0::simple_par_for_impl<migraphx::version_2_14_0::gpu::compile_manager::compile(migraphx::version_2_14_0::module&)::{lambda(auto:1)#1}>(unsigned long, unsigned long, migraphx::version_2_14_0::gpu::compile_manager::compile(migraphx::version_2_14_0::module&)::{lambda(auto:1)#1})::{lambda()#1}::operator()() const::{lambda()#1}>(std::__invoke_other, migraphx::version_2_14_0::simple_par_for_impl<migraphx::version_2_14_0::gpu::compile_manager::compile(migraphx::version_2_14_0::module&)::{lambda(auto:1)#1}>(unsigned long, unsigned long, migraphx::version_2_14_0::gpu::compile_manager::compile(migraphx::version_2_14_0::module&)::{lambda(auto:1)#1})::{lambda()#1}::operator()() const::{lambda()#1}&&) (__f=...) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/invoke.h:61
#32 0x00007421ad886905 in std::__invoke<migraphx::version_2_14_0::simple_par_for_impl<migraphx::version_2_14_0::gpu::compile_manager::compile(migraphx::version_2_14_0::module&)::{lambda(auto:1)#1}>(unsigned long, unsigned long, migraphx::version_2_14_0::gpu::compile_manager::compile(migraphx::version_2_14_0::module&)::{lambda(auto:1)#1})::{lambda()#1}::operator()() const::{lambda()#1}> (__fn=...) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/invoke.h:96
#33 0x00007421ad8868dd in std::thread::_Invoker<std::tuple<migraphx::version_2_14_0::simple_par_for_impl<migraphx::version_2_14_0::gpu::compile_manager::compile(migraphx::version_2_14_0::module&)::{lambda(auto:1)#1}>(unsigned long, unsigned long, migraphx::version_2_14_0::gpu::compile_manager::compile(migraphx::version_2_14_0::module&)::{lambda(auto:1)#1})::{lambda()#1}::operator()() const::{lambda()#1}> >::_M_invoke<0ul>(std::_Index_tuple<0ul>) (this=0x58f1f6eb8508) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/std_thread.h:292
#34 0x00007421ad8868b5 in std::thread::_Invoker<std::tuple<migraphx::version_2_14_0::simple_par_for_impl<migraphx::version_2_14_0::gpu::compile_manager::compile(migraphx::version_2_14_0::module&)::{lambda(auto:1)#1}>(unsigned long, unsigned long, migraphx::version_2_14_0::gpu::compile_manager::compile(migraphx::version_2_14_0::module&)::{lambda(auto:1)#1})::{lambda()#1}::operator()() const::{lambda()#1}> >::operator()() (this=0x58f1f6eb8508) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/std_thread.h:299
#35 0x00007421ad886809 in std::thread::_State_impl<std::thread::_Invoker<std::tuple<migraphx::version_2_14_0::simple_par_for_impl<migraphx::version_2_14_0::gpu::compile_manager::compile(migraphx::version_2_14_0::module&)::{lambda(auto:1)#1}>(unsigned long, unsigned long, migraphx::version_2_14_0::gpu::compile_manager::compile(migraphx::version_2_14_0::module&)::{lambda(auto:1)#1})::{lambda()#1}::operator()() const::{lambda()#1}> > >::_M_run() (this=0x58f1f6eb8500) at /usr/lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/std_thread.h:244
#36 0x00007421807ae67f in ?? () from /usr/lib/x86_64-linux-gnu/libstdc++.so.6
#37 0x0000742180b356db in start_thread (arg=0x7420e1aa6700) at pthread_create.c:463
#38 0x000074217fe3d61f in clone () at ../sysdeps/unix/sysv/linux/x86_64/clone.S:95

The error happens inside hiprtc_src_file::hiprtc_src_file(const src_file& s). It seems src_file::content, which is std::string_view, is invalid. For migraphx/kernels/vectorize.hpp, the data seems fine but size is 123387672153986. So embedding of the file size is somehow failing?

So lets change it:

diff --git a/cmake/Embed.cmake b/cmake/Embed.cmake
index 2e0d055dd..35f929d43 100644
--- a/cmake/Embed.cmake
+++ b/cmake/Embed.cmake
@@ -35,6 +35,7 @@ else()
     endif()
     set_property(CACHE EMBED_USE PROPERTY STRINGS "LD;CArrays")
 endif()
+set(EMBED_USE CArrays)
 
 if(EMBED_USE STREQUAL "LD")
     find_program(EMBED_LD ld REQUIRED)
@@ -216,7 +217,7 @@ function(embed_file FILE BASE_DIRECTORY)
         string(REGEX REPLACE ", $" "" ARRAY_VALUES ${ARRAY_VALUES})
         file(WRITE "${OUTPUT_FILE}" "
 #include <cstddef>
-extern const char _binary_${OUTPUT_SYMBOL}_start[] = { ${ARRAY_VALUES} };
+extern const unsigned char _binary_${OUTPUT_SYMBOL}_start[] = { ${ARRAY_VALUES} };
 extern const size_t _binary_${OUTPUT_SYMBOL}_length = sizeof(_binary_${OUTPUT_SYMBOL}_start);
 ")
         set(OUTPUT_FILE ${OUTPUT_FILE} PARENT_SCOPE)

And it all works now! I can compile stuff. So the LD based embedding fails somehow for my custom TheRock build or linux.

So LD embedding works with this:

diff --git a/cmake/Embed.cmake b/cmake/Embed.cmake
index 2e0d055dd..c50b170a3 100644
--- a/cmake/Embed.cmake
+++ b/cmake/Embed.cmake
@@ -107,12 +107,13 @@ function(generate_embed_source EMBED_NAME EMBED_DIR BASE_DIRECTORY)
             math(EXPR RESOURCE_ID "${RESOURCE_ID} + 1" OUTPUT_FORMAT DECIMAL)
         else()
             set(START_SYMBOL "_binary_${SYMBOL}_start")
+            set(END_SYMBOL "_binary_${SYMBOL}_end")
             set(LENGTH_SYMBOL "_binary_${SYMBOL}_length")
             if(EMBED_USE STREQUAL "LD")
+                set(LENGTH_SYMBOL "static_cast<size_t>(${END_SYMBOL} - ${START_SYMBOL})")
                 string(APPEND EXTERNS "
 extern const char ${START_SYMBOL}[];
-extern const size_t _binary_${SYMBOL}_size;
-const auto ${LENGTH_SYMBOL} = reinterpret_cast<size_t>(&_binary_${SYMBOL}_size);
+extern const char ${END_SYMBOL}[];
 ")
             else()
                 string(APPEND EXTERNS "

But why is the old ld based &_binary_sym_size failing? Apparently because of PIE (position independent executable) and ASLR (address space layout randomization). Some discussion here: https://stackoverflow.com/questions/54844677/why-doesnt-a-linked-binary-files-size-symbol-work-correctly

harakas · 2025-09-29T20:05:14Z

If "config is not required when there is only one kernel" then this PR is not good. Would you like me to modify this PR to fix print_modules() instead and also add a fix for Embed.cmake? @pfultz2

harakas · 2025-09-29T20:07:49Z

Ah yes. The blanket hiding of exceptions. That's quite something. Maybe make it always log out the exception so people can actually see that there are serious problems?

pfultz2 · 2025-09-29T20:16:32Z

Ah yes. The blanket hiding of exceptions. That's quite something.

People complain about the unnecessary noise, because applicability checks do not always have 100% coverage in mlir and are completely missing from CK. So for normal cases it should be ignored.

Maybe make it always log out the exception so people can actually see that there are serious problems?

If all kernels do fail, it should output the exception along with the modules it was compiling, but this might need some more work.

Either way, you can set MIGRAPHX_TRACE_BENCHMARKING=1 to see all the exceptions thrown during this pipeline.

pfultz2 · 2025-09-29T20:17:24Z

Would you like me to modify this PR to fix print_modules() instead and also add a fix for Embed.cmake?

Yes please do.

…_modules

…nt executable (PIE) is enabled

harakas · 2025-10-01T09:57:10Z

I have updated the PR now with check at print_modules and a patch for Embed.cmake. Please review @pfultz2

I also had a cursory look at other ROCm modules in regards to embedding files and I see MIOpen does it properly (https://github.com/ROCm/rocm-libraries/blob/develop/projects/miopen/cmake/embed.cmake#L68-L83), but there is the same problem in CK. Someone should probably fix it: https://github.com/ROCm/composable_kernel/blob/develop/cmake/Embed.cmake#L99-L105

pfultz2 · 2025-10-19T23:52:20Z

I also had a cursory look at other ROCm modules in regards to embedding files and I see MIOpen does it properly (https://github.com/ROCm/rocm-libraries/blob/develop/projects/miopen/cmake/embed.cmake#L68-L83), but there is the same problem in CK. Someone should probably fix it: https://github.com/ROCm/composable_kernel/blob/develop/cmake/Embed.cmake#L99-L105

We used to do it that way but it was changed in #1999.

harakas requested a review from causten as a code owner September 25, 2025 10:47

causten requested a review from Copilot September 28, 2025 16:54

Copilot AI reviewed Sep 28, 2025

View reviewed changes

causten approved these changes Sep 28, 2025

View reviewed changes

pfultz2 reviewed Sep 28, 2025

View reviewed changes

harakas added 2 commits October 1, 2025 12:23

Avoid null dereference when no config provided at compile_plan::print…

a821f6e

…_modules

Fix embedding of files with ld when compiling with position independe…

169a417

…nt executable (PIE) is enabled

harakas force-pushed the better_error branch from afd9720 to 169a417 Compare October 1, 2025 09:42

harakas requested review from causten and pfultz2 October 2, 2025 09:50

Merge branch 'develop' into better_error

5d9a61b

causten approved these changes Oct 17, 2025

View reviewed changes

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

Error check to avoid null config dereference #4319

Error check to avoid null config dereference #4319

harakas commented Sep 25, 2025

Uh oh!

Copilot AI left a comment

Uh oh!

pfultz2 Sep 28, 2025

Uh oh!

pfultz2 commented Sep 28, 2025

Uh oh!

harakas commented Sep 29, 2025 •

edited

Loading

Uh oh!

harakas commented Sep 29, 2025

Uh oh!

harakas commented Sep 29, 2025

Uh oh!

harakas commented Sep 29, 2025

Uh oh!

pfultz2 commented Sep 29, 2025 •

edited

Loading

Uh oh!

pfultz2 commented Sep 29, 2025

Uh oh!

harakas commented Oct 1, 2025

Uh oh!

pfultz2 commented Oct 19, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Error check to avoid null config dereference #4319

Are you sure you want to change the base?

Error check to avoid null config dereference #4319

Conversation

harakas commented Sep 25, 2025

Motivation

Technical Details

Changelog Category

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull Request Overview

Uh oh!

pfultz2 Sep 28, 2025

Choose a reason for hiding this comment

Uh oh!

pfultz2 commented Sep 28, 2025

Uh oh!

harakas commented Sep 29, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

harakas commented Sep 29, 2025

Uh oh!

harakas commented Sep 29, 2025

Uh oh!

harakas commented Sep 29, 2025

Uh oh!

pfultz2 commented Sep 29, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

pfultz2 commented Sep 29, 2025

Uh oh!

harakas commented Oct 1, 2025

Uh oh!

pfultz2 commented Oct 19, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

harakas commented Sep 29, 2025 •

edited

Loading

pfultz2 commented Sep 29, 2025 •

edited

Loading