[GitHub Actions] Changing default behaviour for PR imports by jayhawk-commits · Pull Request #27 · ROCm/rocm-libraries

jayhawk-commits · 2025-05-13T22:48:52Z

No description provided.

* initial attempt to create handle * address review comments and added hip * fix review coments * removed macro * remove hard coded path of hip * remove unused comment * add handle class * compiles now * remove typo * remove comments * remove allocation * format * add try catch * get stream * remove get_stream * some more cleanups * add handle tests * create handle class pointer label in C * add missing copy right * minor change * add name space * fix review comments * use dynamic cast * add more test * minor test fixes * bg/LWPMIOPEN-1473_backend_imp_API_for_handle * Code review comments --------- Co-authored-by: Bibek Ghimire <bghimire@amd.com> Co-authored-by: Brian Harrison <Brian.Harrison@amd.com>

* initial attempt to create handle * address review comments and added hip * fix review coments * removed macro * remove hard coded path of hip * remove unused comment * add handle class * compiles now * remove typo * remove comments * remove allocation * format * add try catch * get stream * remove get_stream * some more cleanups * add handle tests * create handle class pointer label in C * add missing copy right * minor change * add name space * fix review comments * use dynamic cast * add more test * minor test fixes * bg/LWPMIOPEN-1473_backend_imp_API_for_handle * Code review comments --------- Co-authored-by: Bibek Ghimire <bghimire@amd.com> Co-authored-by: Brian Harrison <Brian.Harrison@amd.com> [ROCm/hipDNN commit: afef0c8]

## Motivation Enable gfx1152 and gfx1153. ## Technical Details 1. combine arrays into tables and use local macros to reduce repetition (for maintainability) 2. monkey-see-monkey-do wherever `gfx11...` was found ## Test Plan Build existing ctests for, and run them on, gfx1152 and gfx1153. ## Test Result ### <details> <summary>gfx1152 passed (click to see log)</summary> ``` INFO:root:++ Exec [/tmp/eble]$ ctest --test-dir /tmp/eble/rocm/bin/rocprim --output-o n-failure --parallel 2 --exclude-regex 'rocprim.lookback_reproducibility|rocprim.link ing|rocprim.device_merge_inplace|rocprim.device_merge_sort|rocprim.device_partition|r ocprim.device_radix_sort|rocprim.device_scan|rocprim.device_select|rocprim.device_fin d_first_of|rocprim.device_reduce_by_key' --timeout 60 Test project /tmp/eble/rocm/bin/rocprim Start 1: hip.device_api Start 2: hip.async_copy 1/73 Test #2: hip.async_copy .............................. Passed 0.01 sec Start 3: hip.ordered_block_id 2/73 Test #1: hip.device_api .............................. Passed 0.02 sec Start 4: rocprim.internal_merge_path 3/73 Test #4: rocprim.internal_merge_path ................. Passed 0.01 sec Start 5: rocprim.basic_test 4/73 Test #3: hip.ordered_block_id ........................ Passed 0.01 sec Start 6: rocprim.arg_index_iterator 5/73 Test #5: rocprim.basic_test .......................... Passed 0.01 sec Start 7: rocprim.temporary_storage_partitioning 6/73 Test #6: rocprim.arg_index_iterator .................. Passed 0.01 sec Start 8: rocprim.block_adjacent_difference 7/73 Test #7: rocprim.temporary_storage_partitioning ...... Passed 0.01 sec Start 9: rocprim.block_discontinuity 8/73 Test #8: rocprim.block_adjacent_difference ........... Passed 2.34 sec Start 10: rocprim.bit_cast 9/73 Test #10: rocprim.bit_cast ............................ Passed 0.02 sec Start 11: rocprim.block_exchange 10/73 Test #11: rocprim.block_exchange ...................... Passed 0.73 sec Start 12: rocprim.block_histogram 11/73 Test #12: rocprim.block_histogram ..................... Passed 0.54 sec Start 13: rocprim.block_load_store 12/73 Test #13: rocprim.block_load_store .................... Passed 0.44 sec Start 14: rocprim.block_sort_merge 13/73 Test #14: rocprim.block_sort_merge .................... Passed 0.02 sec Start 15: rocprim.block_sort_merge_stable 14/73 Test #15: rocprim.block_sort_merge_stable ............. Passed 0.02 sec Start 16: rocprim.block_radix_rank 15/73 Test #16: rocprim.block_radix_rank .................... Passed 0.03 sec Start 17: rocprim.block_radix_sort 16/73 Test #17: rocprim.block_radix_sort .................... Passed 4.79 sec Start 18: rocprim.block_reduce 17/73 Test #18: rocprim.block_reduce ........................ Passed 0.26 sec Start 19: rocprim.block_run_length_decode 18/73 Test #19: rocprim.block_run_length_decode ............. Passed 0.54 sec Start 20: rocprim.block_scan 19/73 Test #20: rocprim.block_scan .......................... Passed 0.04 sec Start 21: rocprim.block_shuffle 20/73 Test #21: rocprim.block_shuffle ....................... Passed 2.70 sec Start 22: rocprim.block_sort_bitonic 21/73 Test #9: rocprim.block_discontinuity ................. Passed 17.36 sec Start 23: rocprim.config_dispatch 22/73 Test #23: rocprim.config_dispatch ..................... Passed 0.09 sec Start 24: rocprim.constant_iterator 23/73 Test #24: rocprim.constant_iterator ................... Passed 0.07 sec Start 25: rocprim.counting_iterator 24/73 Test #25: rocprim.counting_iterator ................... Passed 0.07 sec Start 26: rocprim.device_batch_memcpy 25/73 Test #26: rocprim.device_batch_memcpy ................. Passed 1.20 sec Start 27: rocprim.device_binary_search 26/73 Test #27: rocprim.device_binary_search ................ Passed 0.02 sec Start 28: rocprim.device_adjacent_difference 27/73 Test #28: rocprim.device_adjacent_difference .......... Passed 0.01 sec Start 29: rocprim.device_adjacent_find 28/73 Test #29: rocprim.device_adjacent_find ................ Passed 0.01 sec Start 30: rocprim.device_find_end 29/73 Test #30: rocprim.device_find_end ..................... Passed 0.01 sec Start 31: rocprim.device_histogram 30/73 Test #22: rocprim.block_sort_bitonic .................. Passed 13.23 sec Start 32: rocprim.device_merge 31/73 Test #31: rocprim.device_histogram .................... Passed 7.85 sec Start 33: rocprim.nth_element 32/73 Test #33: rocprim.nth_element ......................... Passed 0.03 sec Start 34: rocprim.device_partial_sort 33/73 Test #34: rocprim.device_partial_sort ................. Passed 0.02 sec Start 35: rocprim.device_reduce 34/73 Test #35: rocprim.device_reduce ....................... Passed 8.94 sec Start 36: rocprim.device_run_length_encode 35/73 Test #32: rocprim.device_merge ........................ Passed 14.05 sec Start 37: rocprim.device_search 36/73 Test #37: rocprim.device_search ....................... Passed 0.02 sec Start 38: rocprim.device_segmented_radix_sort 37/73 Test #36: rocprim.device_run_length_encode ............ Passed 13.92 sec Start 39: rocprim.device_search_n 38/73 Test #39: rocprim.device_search_n ..................... Passed 0.02 sec Start 40: rocprim.device_segmented_reduce 39/73 Test #40: rocprim.device_segmented_reduce ............. Passed 5.21 sec Start 41: rocprim.device_segmented_scan 40/73 Test #41: rocprim.device_segmented_scan ............... Passed 0.02 sec Start 42: rocprim.device_transform 41/73 Test #42: rocprim.device_transform .................... Passed 13.90 sec Start 43: rocprim.discard_iterator 42/73 Test #43: rocprim.discard_iterator .................... Passed 0.07 sec Start 44: rocprim.radix_key_codec 43/73 Test #44: rocprim.radix_key_codec ..................... Pas09:54:11 [55/1943] Start 45: rocprim.predicate_iterator 44/73 Test #45: rocprim.predicate_iterator .................. Passed 0.07 sec Start 46: rocprim.reverse_iterator 45/73 Test #46: rocprim.reverse_iterator .................... Passed 0.09 sec Start 47: rocprim.rocprim_tuple 46/73 Test #47: rocprim.rocprim_tuple ....................... Passed 0.01 sec Start 48: rocprim.rocprim_types 47/73 Test #48: rocprim.rocprim_types ....................... Passed 0.01 sec Start 49: rocprim.texture_cache_iterator 48/73 Test #49: rocprim.texture_cache_iterator .............. Passed 0.01 sec Start 50: rocprim.thread 49/73 Test #50: rocprim.thread .............................. Passed 0.07 sec Start 51: rocprim.thread_algos 50/73 Test #51: rocprim.thread_algos ........................ Passed 0.35 sec Start 52: rocprim.tuple 51/73 Test #52: rocprim.tuple ............................... Passed 0.02 sec Start 53: rocprim.utils_sort_checker 52/73 Test #53: rocprim.utils_sort_checker .................. Passed 0.01 sec Start 54: rocprim.transform_iterator 53/73 Test #54: rocprim.transform_iterator .................. Passed 0.11 sec Start 55: rocprim.type_traits_interface_cpp17 54/73 Test #55: rocprim.type_traits_interface_cpp17 ......... Passed 0.01 sec Start 56: rocprim.type_traits_interface_gnupp17 55/73 Test #56: rocprim.type_traits_interface_gnupp17 ....... Passed 0.01 sec Start 57: rocprim.type_traits_interface_cpp20 56/73 Test #57: rocprim.type_traits_interface_cpp20 ......... Passed 0.01 sec Start 58: rocprim.type_traits_interface_gnupp20 57/73 Test #58: rocprim.type_traits_interface_gnupp20 ....... Passed 0.01 sec Start 59: rocprim.no_half_operators 58/73 Test #59: rocprim.no_half_operators ................... Passed 0.01 sec Start 60: rocprim.intrinsics 59/73 Test #60: rocprim.intrinsics .......................... Passed 0.21 sec Start 61: rocprim.intrinsics_atomic 60/73 Test #61: rocprim.intrinsics_atomic ................... Passed 0.02 sec Start 62: rocprim.invoke_result 61/73 Test #62: rocprim.invoke_result ....................... Passed 0.01 sec Start 63: rocprim.warp_exchange 62/73 Test #63: rocprim.warp_exchange ....................... Passed 0.08 sec Start 64: rocprim.warp_load 63/73 Test #64: rocprim.warp_load ........................... Passed 0.08 sec Start 65: rocprim.warp_reduce 64/73 Test #65: rocprim.warp_reduce ......................... Passed 0.14 sec Start 66: rocprim.warp_scan 65/73 Test #66: rocprim.warp_scan ........................... Passed 0.20 sec Start 67: rocprim.warp_scan_disable_dpp_disable_dpp 66/73 Test #67: rocprim.warp_scan_disable_dpp_disable_dpp ... Passed 0.21 sec Start 68: rocprim.warp_sort 67/73 Test #68: rocprim.warp_sort ........................... Passed 0.09 sec Start 69: rocprim.warp_store 68/73 Test #69: rocprim.warp_store .......................... Passed 0.02 sec Start 70: rocprim.zip_iterator 69/73 Test #70: rocprim.zip_iterator ........................ Passed 0.02 sec Start 71: rocprim.accumulator_t 70/73 Test #71: rocprim.accumulator_t ....................... Passed 0.02 sec Start 72: hipgraph.basic 71/73 Test #72: hipgraph.basic .............................. Passed 0.02 sec Start 73: hipgraph.algs 72/73 Test #73: hipgraph.algs ............................... Passed 0.01 sec 73/73 Test #38: rocprim.device_segmented_radix_sort ......... Passed 31.97 sec 100% tests passed, 0 tests failed out of 73 Total Test time (real) = 71.80 sec ✅ test_rocprim.py PASSED ``` </details> <details> <summary>gfx1153 passed (click to see log)</summary> ``` INFO:root:++ Exec [/tmp/eble]$ ctest --test-dir /tmp/eble/rocm/bin/rocprim --output-o n-failure --parallel 1 --exclude-regex 'rocprim.lookback_reproducibility|rocprim.link ing|rocprim.device_merge_inplace|rocprim.device_merge_sort|rocprim.device_partition|r ocprim.device_radix_sort|rocprim.device_scan|rocprim.device_select|rocprim.device_fin d_first_of|rocprim.device_reduce_by_key' --timeout 60 Test project /tmp/eble/rocm/bin/rocprim Start 1: hip.device_api 1/73 Test #1: hip.device_api .............................. Passed 0.01 sec Start 2: hip.async_copy 2/73 Test #2: hip.async_copy .............................. Passed 0.01 sec Start 3: hip.ordered_block_id 3/73 Test #3: hip.ordered_block_id ........................ Passed 0.01 sec Start 4: rocprim.internal_merge_path 4/73 Test #4: rocprim.internal_merge_path ................. Passed 0.01 sec Start 5: rocprim.basic_test 5/73 Test #5: rocprim.basic_test .......................... Passed 0.01 sec Start 6: rocprim.arg_index_iterator 6/73 Test #6: rocprim.arg_index_iterator .................. Passed 0.01 sec Start 7: rocprim.temporary_storage_partitioning 7/73 Test #7: rocprim.temporary_storage_partitioning ...... Passed 0.01 sec Start 8: rocprim.block_adjacent_difference 8/73 Test #8: rocprim.block_adjacent_difference ........... Passed 2.94 sec Start 9: rocprim.block_discontinuity 9/73 Test #9: rocprim.block_discontinuity ................. Passed 21.10 sec Start 10: rocprim.bit_cast 10/73 Test #10: rocprim.bit_cast ............................ Passed 0.01 sec Start 11: rocprim.block_exchange 11/73 Test #11: rocprim.block_exchange ...................... Passed 2.20 sec Start 12: rocprim.block_histogram 12/73 Test #12: rocprim.block_histogram ..................... Passed 0.71 sec Start 13: rocprim.block_load_store 13/73 Test #13: rocprim.block_load_store .................... Passed 0.48 sec Start 14: rocprim.block_sort_merge 14/73 Test #14: rocprim.block_sort_merge .................... Passed 0.02 sec Start 15: rocprim.block_sort_merge_stable 15/73 Test #15: rocprim.block_sort_merge_stable ............. Passed 0.02 sec Start 16: rocprim.block_radix_rank 16/73 Test #16: rocprim.block_radix_rank .................... Passed 0.02 sec Start 17: rocprim.block_radix_sort 17/73 Test #17: rocprim.block_radix_sort .................... Passed 6.12 sec Start 18: rocprim.block_reduce 18/73 Test #18: rocprim.block_reduce ........................ Passed 0.31 sec Start 19: rocprim.block_run_length_decode 19/73 Test #19: rocprim.block_run_length_decode ............. Passed 0.68 sec Start 20: rocprim.block_scan 20/73 Test #20: rocprim.block_scan .......................... Passed 0.03 sec Start 21: rocprim.block_shuffle 21/73 Test #21: rocprim.block_shuffle ....................... Passed 3.63 sec Start 22: rocprim.block_sort_bitonic 22/73 Test #22: rocprim.block_sort_bitonic .................. Passed 19.34 sec Start 23: rocprim.config_dispatch 23/73 Test #23: rocprim.config_dispatch ..................... Passed 0.10 sec Start 24: rocprim.constant_iterator 24/73 Test #24: rocprim.constant_iterator ................... Passed 0.09 sec Start 25: rocprim.counting_iterator 25/73 Test #25: rocprim.counting_iterator ................... Passed 0.09 sec Start 26: rocprim.device_batch_memcpy 26/73 Test #26: rocprim.device_batch_memcpy ................. Passed 1.42 sec Start 27: rocprim.device_binary_search 27/73 Test #27: rocprim.device_binary_search ................ Passed 0.01 sec Start 28: rocprim.device_adjacent_difference 28/73 Test #28: rocprim.device_adjacent_difference .......... Passed 0.01 sec Start 29: rocprim.device_adjacent_find 29/73 Test #29: rocprim.device_adjacent_find ................ Passed 0.01 sec Start 30: rocprim.device_find_end 30/73 Test #30: rocprim.device_find_end ..................... Passed 0.01 sec Start 31: rocprim.device_histogram 31/73 Test #31: rocprim.device_histogram .................... Passed 8.77 sec Start 32: rocprim.device_merge 32/73 Test #32: rocprim.device_merge ........................ Passed 16.23 sec Start 33: rocprim.nth_element 33/73 Test #33: rocprim.nth_element ......................... Passed 0.01 sec Start 34: rocprim.device_partial_sort 34/73 Test #34: rocprim.device_partial_sort ................. Passed 0.02 sec Start 35: rocprim.device_reduce 35/73 Test #35: rocprim.device_reduce ....................... Passed 10.92 sec Start 36: rocprim.device_run_length_encode 36/73 Test #36: rocprim.device_run_length_encode ............ Passed 14.34 sec Start 37: rocprim.device_search 37/73 Test #37: rocprim.device_search ....................... Passed 0.01 sec Start 38: rocprim.device_segmented_radix_sort 38/73 Test #38: rocprim.device_segmented_radix_sort ......... Passed 38.28 sec Start 39: rocprim.device_search_n 39/73 Test #39: rocprim.device_search_n ..................... Passed 0.02 sec Start 40: rocprim.device_segmented_reduce 40/73 Test #40: rocprim.device_segmented_reduce ............. Passed 7.19 sec Start 41: rocprim.device_segmented_scan 41/73 Test #41: rocprim.device_segmented_scan ............... Passed 0.02 sec Start 42: rocprim.device_transform 42/73 Test #42: rocprim.device_transform .................... Passed 17.64 sec Start 43: rocprim.discard_iterator 43/73 Test #43: rocprim.discard_iterator .................... Passed 0.12 sec Start 44: rocprim.radix_key_codec 44/73 Test #44: rocprim.radix_key_codec ..................... Passed 0.01 sec Start 45: rocprim.predicate_iterator 45/73 Test #45: rocprim.predicate_iterator .................. Passed 0.08 sec Start 46: rocprim.reverse_iterator 46/73 Test #46: rocprim.reverse_iterator .................... Pas10:13:26 [46/1844] Start 47: rocprim.rocprim_tuple 47/73 Test #47: rocprim.rocprim_tuple ....................... Passed 0.01 sec Start 48: rocprim.rocprim_types 48/73 Test #48: rocprim.rocprim_types ....................... Passed 0.01 sec Start 49: rocprim.texture_cache_iterator 49/73 Test #49: rocprim.texture_cache_iterator .............. Passed 0.01 sec Start 50: rocprim.thread 50/73 Test #50: rocprim.thread .............................. Passed 0.08 sec Start 51: rocprim.thread_algos 51/73 Test #51: rocprim.thread_algos ........................ Passed 0.43 sec Start 52: rocprim.tuple 52/73 Test #52: rocprim.tuple ............................... Passed 0.01 sec Start 53: rocprim.utils_sort_checker 53/73 Test #53: rocprim.utils_sort_checker .................. Passed 0.01 sec Start 54: rocprim.transform_iterator 54/73 Test #54: rocprim.transform_iterator .................. Passed 0.12 sec Start 55: rocprim.type_traits_interface_cpp17 55/73 Test #55: rocprim.type_traits_interface_cpp17 ......... Passed 0.01 sec Start 56: rocprim.type_traits_interface_gnupp17 56/73 Test #56: rocprim.type_traits_interface_gnupp17 ....... Passed 0.01 sec Start 57: rocprim.type_traits_interface_cpp20 57/73 Test #57: rocprim.type_traits_interface_cpp20 ......... Passed 0.01 sec Start 58: rocprim.type_traits_interface_gnupp20 58/73 Test #58: rocprim.type_traits_interface_gnupp20 ....... Passed 0.01 sec Start 59: rocprim.no_half_operators 59/73 Test #59: rocprim.no_half_operators ................... Passed 0.01 sec Start 60: rocprim.intrinsics 60/73 Test #60: rocprim.intrinsics .......................... Passed 0.29 sec Start 61: rocprim.intrinsics_atomic 61/73 Test #61: rocprim.intrinsics_atomic ................... Pas10:13:27 [16/1844] Start 62: rocprim.invoke_result 62/73 Test #62: rocprim.invoke_result ....................... Passed 0.01 sec Start 63: rocprim.warp_exchange 63/73 Test #63: rocprim.warp_exchange ....................... Passed 0.09 sec Start 64: rocprim.warp_load 64/73 Test #64: rocprim.warp_load ........................... Passed 0.09 sec Start 65: rocprim.warp_reduce 65/73 Test #65: rocprim.warp_reduce ......................... Passed 0.17 sec Start 66: rocprim.warp_scan 66/73 Test #66: rocprim.warp_scan ........................... Passed 0.26 sec Start 67: rocprim.warp_scan_disable_dpp_disable_dpp 67/73 Test #67: rocprim.warp_scan_disable_dpp_disable_dpp ... Passed 0.26 sec Start 68: rocprim.warp_sort 68/73 Test #68: rocprim.warp_sort ........................... Passed 0.10 sec Start 69: rocprim.warp_store 69/73 Test #69: rocprim.warp_store .......................... Passed 0.01 sec Start 70: rocprim.zip_iterator 70/73 Test #70: rocprim.zip_iterator ........................ Passed 0.01 sec Start 71: rocprim.accumulator_t 71/73 Test #71: rocprim.accumulator_t ....................... Passed 0.01 sec Start 72: hipgraph.basic 72/73 Test #72: hipgraph.basic .............................. Passed 0.01 sec Start 73: hipgraph.algs 73/73 Test #73: hipgraph.algs ............................... Passed 0.01 sec 100% tests passed, 0 tests failed out of 73 Total Test time (real) = 175.34 sec ✅ test_rocprim.py PASSED ``` <detail> ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

Enable wmma-related tests on gfx1250

…being used after being freed. (#5220) ## Motivation A `heap-use-after-free` error was triggered by AddressSanitizer on test `CPU_Dump_NAN_FP32.testDump`. ## Technical Details Root Cause Analysis: The AddressSanitizer error occurred because the HIPOCProgramImpl constructor was not storing the binary data passed to it. When LoadProgram called LoadBinary and created a HIPOCProgram with the returned vector, the temporary vector would go out of scope, but COMGR still needed to access the binary data later, causing a use-after-free. - The fix ensures that the HIPOCProgramImpl object owns the binary data for its entire lifetime - Both constructors now consistently store the binary data in the `binary` member variable (std::vector) - The uint8_t constructor converts the data to char format using iterator range construction - This prevents the use-after-free that occurred when COMGR tried to access freed memory ## Test Plan Test output before change: ``` HSA_XNACK=1 ASAN_OPTIONS=symbolize=1 ./build/ml-libs/MIOpen/build/bin/miopen_gtest --gtest_filter="*CPU_Dump_NAN_FP32*" PRNG seed: 12345678 Note: Google Test filter = *CPU_Dump_NAN_FP32* [==========] Running 1 test from 1 test suite. [----------] Global test environment set-up. [----------] 1 test from CPU_Dump_NAN_FP32 [ RUN ] CPU_Dump_NAN_FP32.testDump ================================================================= ==3639==ERROR: AddressSanitizer: heap-use-after-free on address 0x7e0f08c50200 at pc 0x7f5f8d7a6554 bp 0x7ffcb7c4a730 sp 0x7ffcb7c49ee8 READ of size 26088 at 0x7e0f08c50200 thread T0 #0 0x7f5f8d7a6553 in memcpy /data/nhanna/repos/TheRock/compiler/amd-llvm/compiler-rt/lib/asan/../sanitizer_common/sanitizer_common_interceptors_memintrinsics.inc:117:5 #1 0x7f5f23d61d78 in COMGR::setCStr(char*&, llvm::StringRef, unsigned long*) /data/nhanna/repos/TheRock/compiler/amd-llvm/amd/comgr/src/comgr.cpp:216:9 #2 0x7f5f23d61d78 in COMGR::DataObject::setData(llvm::StringRef) /data/nhanna/repos/TheRock/compiler/amd-llvm/amd/comgr/src/comgr.cpp:334:17 #3 0x7f5f23d61d78 in amd_comgr_set_data /data/nhanna/repos/TheRock/compiler/amd-llvm/amd/comgr/src/comgr.cpp:606:24 #4 0x7f5f221dc1d3 in amd::Comgr::set_data(amd_comgr_data_s, unsigned long, char const*) /data/nhanna/repos/TheRock/rocm-systems/projects/clr/rocclr/device/comgrctx.hpp:252:12 #5 0x7f5f221dc1d3 in amd::device::Program::getSymbolsFromCodeObj(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>>>>*, amd_comgr_symbol_type_s) const /data/nhanna/repos/TheRock/rocm-systems/projects/clr/rocclr/device/devprogram.cpp:2061:14 #6 0x7f5f219e6f7c in hip::DynCO::populateDynGlobalVars() /data/nhanna/repos/TheRock/rocm-systems/projects/clr/hipamd/src/hip_code_object.cpp:216:22 #7 0x7f5f219e8e6a in hip::DynCO::getDynFunc(ihipModuleSymbol_t**, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>) /data/nhanna/repos/TheRock/rocm-systems/projects/clr/hipamd/src/hip_code_object.cpp:125:22 #8 0x7f5f21f842ba in hip::PlatformState::GetDynFunc(ihipModuleSymbol_t**, ihipModule_t*, char const*) /data/nhanna/repos/TheRock/rocm-systems/projects/clr/hipamd/src/hip_platform.cpp:884:22 #9 0x7f5f21ec2d71 in hip::hipModuleGetFunction(ihipModuleSymbol_t**, ihipModule_t*, char const*) /data/nhanna/repos/TheRock/rocm-systems/projects/clr/hipamd/src/hip_module.cpp:89:47 #10 0x7f5f2212c588 in hipModuleGetFunction /data/nhanna/repos/TheRock/rocm-systems/projects/clr/hipamd/src/hip_table_interface.cpp:1926:10 #11 0x7f5f7478d806 in miopen::HIPOCKernel::HIPOCKernel(miopen::HIPOCProgram, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>, std::vector<unsigned long, std::allocator<unsigned long>>, std::vector<unsigned long, std::allocator<unsigned long>>) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/include/miopen/hipoc_kernel.hpp:225:25 #12 0x7f5f766febb7 in miopen::KernelCache::AddKernel(miopen::Handle const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>, unsigned long, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, miopen::HIPOCProgram*) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/kernel_cache.cpp:161:18 #13 0x7f5f76b6f0e4 in miopen::Handle::AddKernel(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, unsigned long, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&) const /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/hip/handlehip.cpp:450:34 #14 0x7f5f7411b52f in miopen::checkNumericsImpl(miopen::Handle const&, int, miopen::TensorDescriptor const&, void const*, bool) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/check_numerics.cpp:107:12 #15 0x55e87c72ebee in void testDumpWithNan<float>(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/test/gtest/dumpTensorTest.cpp:130:8 #16 0x55e87c72d4e8 in CPU_Dump_NAN_FP32_testDump_Test::TestBody() /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/test/gtest/dumpTensorTest.cpp:157:37 #17 0x55e87ef19d5e in void testing::internal::HandleSehExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2653:27 #18 0x55e87ef19d5e in void testing::internal::HandleExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2689:52 #19 0x55e87ef04cdd in testing::Test::Run() /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2728:50 #20 0x55e87ef04cdd in testing::Test::Run() /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2718:6 #21 0x55e87ef04e64 in testing::TestInfo::Run() /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2874:14 #22 0x55e87ef0500e in testing::TestSuite::Run() /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:3052:33 #23 0x55e87ef0500e in testing::TestSuite::Run() /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:3006:6 #24 0x55e87ef0d20b in testing::internal::UnitTestImpl::RunAllTests() /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:6004:47 #25 0x55e87ef1a1de in bool testing::internal::HandleSehExceptionsInMethodIfSupported<testing::internal::UnitTestImpl, bool>(testing::internal::UnitTestImpl*, bool (testing::internal::UnitTestImpl::*)(), char const*) /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2653:27 #26 0x55e87ef1a1de in bool testing::internal::HandleExceptionsInMethodIfSupported<testing::internal::UnitTestImpl, bool>(testing::internal::UnitTestImpl*, bool (testing::internal::UnitTestImpl::*)(), char const*) /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2689:52 #27 0x55e87ef051b5 in testing::UnitTest::Run() /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:5583:55 #28 0x55e87eee3f9b in RUN_ALL_TESTS() /data/nhanna/repos/TheRock/build/third-party/googletest/dist/include/gtest/gtest.h:2334:73 #29 0x55e87eee3f9b in main /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/test/gtest/main_hip.cpp:34:12 #30 0x7f5f20a587e4 in __libc_start_main (/lib64/libc.so.6+0x3a7e4) (BuildId: 889235a2805b8308b2d0274921bbe1890e9a1986) #31 0x55e87b0bcf2d in _start (/data/nhanna/repos/TheRock/build/ml-libs/MIOpen/build/bin/miopen_gtest+0x126bf2d) 0x7e0f08c50200 is located 0 bytes inside of 26088-byte region [0x7e0f08c50200,0x7e0f08c567e8) freed by thread T0 here: #0 0x7f5f8d7b8ba2 in operator delete(void*, unsigned long) /data/nhanna/repos/TheRock/compiler/amd-llvm/compiler-rt/lib/asan/asan_new_delete.cpp:190:3 #1 0x7f5f76b7317d in std::__new_allocator<char>::deallocate(char*, unsigned long) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/new_allocator.h:172:2 #2 0x7f5f76b7317d in std::allocator<char>::deallocate(char*, unsigned long) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/allocator.h:210:25 #3 0x7f5f76b7317d in std::allocator_traits<std::allocator<char>>::deallocate(std::allocator<char>&, char*, unsigned long) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/alloc_traits.h:517:13 #4 0x7f5f76b7317d in std::_Vector_base<char, std::allocator<char>>::_M_deallocate(char*, unsigned long) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/stl_vector.h:390:4 #5 0x7f5f76b7317d in std::_Vector_base<char, std::allocator<char>>::~_Vector_base() /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/stl_vector.h:369:2 #6 0x7f5f76b7317d in std::vector<char, std::allocator<char>>::~vector() /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/stl_vector.h:738:7 #7 0x7f5f76b7317d in miopen::Handle::LoadProgram(std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, bool) const /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/hip/handlehip.cpp:633:5 #8 0x7f5f766fda82 in miopen::KernelCache::AddKernel(miopen::Handle const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>, unsigned long, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, miopen::HIPOCProgram*)::'lambda'()::operator()() const /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/kernel_cache.cpp:143:30 #9 0x7f5f766fda82 in miopen::KernelCache::AddKernel(miopen::Handle const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>, unsigned long, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, miopen::HIPOCProgram*) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/kernel_cache.cpp:124:26 #10 0x7f5f76b6f0e4 in miopen::Handle::AddKernel(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, unsigned long, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&) const /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/hip/handlehip.cpp:450:34 #11 0x7f5f7411b52f in miopen::checkNumericsImpl(miopen::Handle const&, int, miopen::TensorDescriptor const&, void const*, bool) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/check_numerics.cpp:107:12 #12 0x55e87c72ebee in void testDumpWithNan<float>(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/test/gtest/dumpTensorTest.cpp:130:8 #13 0x55e87c72d4e8 in CPU_Dump_NAN_FP32_testDump_Test::TestBody() /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/test/gtest/dumpTensorTest.cpp:157:37 #14 0x55e87ef19d5e in void testing::internal::HandleSehExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2653:27 #15 0x55e87ef19d5e in void testing::internal::HandleExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2689:52 previously allocated by thread T0 here: #0 0x7f5f8d7b7f9d in operator new(unsigned long) /data/nhanna/repos/TheRock/compiler/amd-llvm/compiler-rt/lib/asan/asan_new_delete.cpp:109:35 #1 0x7f5f76b720a5 in std::__new_allocator<char>::allocate(unsigned long, void const*) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/new_allocator.h:151:27 #2 0x7f5f76b720a5 in std::allocator<char>::allocate(unsigned long) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/allocator.h:198:32 #3 0x7f5f76b720a5 in std::allocator_traits<std::allocator<char>>::allocate(std::allocator<char>&, unsigned long) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/alloc_traits.h:482:20 #4 0x7f5f76b720a5 in std::_Vector_base<char, std::allocator<char>>::_M_allocate(unsigned long) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/stl_vector.h:381:20 #5 0x7f5f76b720a5 in std::_Vector_base<char, std::allocator<char>>::_M_create_storage(unsigned long) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/stl_vector.h:398:33 #6 0x7f5f76b720a5 in std::_Vector_base<char, std::allocator<char>>::_Vector_base(unsigned long, std::allocator<char> const&) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/stl_vector.h:335:9 #7 0x7f5f76b720a5 in std::vector<char, std::allocator<char>>::vector(std::vector<char, std::allocator<char>> const&) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/stl_vector.h:602:9 #8 0x7f5f76b720a5 in miopen::Handle::LoadProgram(std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, bool) const /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/hip/handlehip.cpp:623:27 #9 0x7f5f766fda82 in miopen::KernelCache::AddKernel(miopen::Handle const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>, unsigned long, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, miopen::HIPOCProgram*)::'lambda'()::operator()() const /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/kernel_cache.cpp:143:30 #10 0x7f5f766fda82 in miopen::KernelCache::AddKernel(miopen::Handle const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>, unsigned long, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, miopen::HIPOCProgram*) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/kernel_cache.cpp:124:26 #11 0x7f5f76b6f0e4 in miopen::Handle::AddKernel(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, unsigned long, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&) const /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/hip/handlehip.cpp:450:34 #12 0x7f5f7411b52f in miopen::checkNumericsImpl(miopen::Handle const&, int, miopen::TensorDescriptor const&, void const*, bool) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/check_numerics.cpp:107:12 #13 0x55e87c72ebee in void testDumpWithNan<float>(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/test/gtest/dumpTensorTest.cpp:130:8 #14 0x55e87c72d4e8 in CPU_Dump_NAN_FP32_testDump_Test::TestBody() /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/test/gtest/dumpTensorTest.cpp:157:37 #15 0x55e87ef19d5e in void testing::internal::HandleSehExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2653:27 #16 0x55e87ef19d5e in void testing::internal::HandleExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2689:52 SUMMARY: AddressSanitizer: heap-use-after-free /data/nhanna/repos/TheRock/compiler/amd-llvm/amd/comgr/src/comgr.cpp:216:9 in COMGR::setCStr(char*&, llvm::StringRef, unsigned long*) Shadow bytes around the buggy address: 0x7e0f08c4ff80: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa 0x7e0f08c50000: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa 0x7e0f08c50080: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa 0x7e0f08c50100: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa 0x7e0f08c50180: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa =>0x7e0f08c50200:[fd]fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd 0x7e0f08c50280: fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd 0x7e0f08c50300: fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd 0x7e0f08c50380: fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd 0x7e0f08c50400: fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd 0x7e0f08c50480: fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd Shadow byte legend (one shadow byte represents 8 application bytes): Addressable: 00 Partially addressable: 01 02 03 04 05 06 07 Heap left redzone: fa Freed heap region: fd Stack left redzone: f1 Stack mid redzone: f2 Stack right redzone: f3 Stack after return: f5 Stack use after scope: f8 Global redzone: f9 Global init order: f6 Poisoned by user: f7 Container overflow: fc Array cookie: ac Intra object redzone: bb ASan internal: fe Left alloca redzone: ca Right alloca redzone: cb ==3639==ABORTING ``` ## Test Result Test output after change: ``` HSA_XNACK=1 ASAN_OPTIONS=symbolize=1 ./build/ml-libs/MIOpen/build/bin/miopen_gtest --gtest_filter="*CPU_Dump_NAN_FP32*" PRNG seed: 12345678 Note: Google Test filter = *CPU_Dump_NAN_FP32* [==========] Running 1 test from 1 test suite. [----------] Global test environment set-up. [----------] 1 test from CPU_Dump_NAN_FP32 [ RUN ] CPU_Dump_NAN_FP32.testDump [ OK ] CPU_Dump_NAN_FP32.testDump (51 ms) [----------] 1 test from CPU_Dump_NAN_FP32 (51 ms total) [----------] Global test environment tear-down [==========] 1 test from 1 test suite ran. (52 ms total) [ PASSED ] 1 test. ``` ## Cline Analysis ### Test Coverage Analysis: __1. LoadProgram Code Path (std::vector constructor):__ - __Primary Test__: `rocm-libraries/projects/miopen/test/gtest/db_sync.cpp` - __Function__: `BuildKernel()` calls `handle.LoadProgram(program_file, program_args, "")` - __Coverage__: This test extensively exercises the LoadProgram → LoadBinary → HIPOCProgramImpl constructor path - __Scope__: Tests multiple GPU architectures (gfx908, gfx90a, gfx942, gfx1030) with different CU counts - __Frequency__: Runs on thousands of kernel configurations in the database sync tests __2. Solution Binary Serialization (std::vector usage):__ - __Primary Test__: `rocm-libraries/projects/miopen/test/gtest/find_2_conv.cpp` - __Function__: `miopenSaveSolution()` and `miopenLoadSolution()` with `std::vector<char> solution_binary` - __Coverage__: Tests the save/load cycle of solution binaries - __Scope__: Tests all convolution directions (Forward, BackwardData, BackwardWeights) __3. Additional Coverage:__ - __Cache Tests__: `rocm-libraries/projects/miopen/test/gtest/cache.cpp` tests compression/decompression with `std::vector<char>` - __Dropout Tests__: Uses `std::vector<unsigned char>` for reserve space (related pattern) __Test Quality Assessment:__ ✅ __Both constructors are well-tested__: - The `std::vector<char>` constructor is heavily exercised through database sync tests - The `std::vector<uint8_t>` constructor would be tested through any code paths that use uint8_t binary data ✅ __Real-world scenarios covered__: - Database synchronization (production kernel loading) - Solution serialization (runtime binary handling) - Multi-threaded execution (db_sync uses up to 32 threads) ✅ __Comprehensive architecture coverage__: - Tests run on multiple GPU architectures - Different compute unit configurations tested __Confidence Level__: Very High ### Performance Analysis: Regarding the performance impact of this fix, it's actually quite minimal and represents good engineering practice: __Memory Impact:__ - __Additional Memory Usage__: Each HIPOCProgramImpl object now stores a copy of the binary data in its `binary` member variable - __Typical Size__: GPU code objects are usually relatively small (typically a few KB to a few MB depending on kernel complexity) - __Lifetime__: The memory is only held for the lifetime of the HIPOCProgram object, which is typically short-lived during kernel loading __Performance Characteristics:__ - __One-time Copy Cost__: There's a single memory copy operation during construction (std::vector copy or iterator range construction) - __No Runtime Overhead__: Once constructed, there's no additional performance cost during kernel execution - __Memory Safety Benefit__: Eliminates potential crashes and undefined behavior, which far outweighs the small memory cost __Context in MIOpen:__ - This occurs during the kernel loading phase, not during actual ML inference/training - Kernel loading is already an expensive operation involving compilation, module creation, etc. - The additional memory copy is negligible compared to the overall kernel loading time __Trade-off Analysis:__ - __Cost__: Small increase in memory usage during kernel loading - __Benefit__: Eliminates memory safety bugs that could cause crashes or data corruption - __Net Result__: Significantly positive - reliability and correctness are much more valuable than the minimal memory overhead In practice, this fix follows the RAII (Resource Acquisition Is Initialization) principle and ensures proper ownership semantics, which is standard best practice in modern C++. The performance impact should be unnoticeable in real-world usage.

…being used after being freed. (ROCm#5220) ## Motivation A `heap-use-after-free` error was triggered by AddressSanitizer on test `CPU_Dump_NAN_FP32.testDump`. ## Technical Details Root Cause Analysis: The AddressSanitizer error occurred because the HIPOCProgramImpl constructor was not storing the binary data passed to it. When LoadProgram called LoadBinary and created a HIPOCProgram with the returned vector, the temporary vector would go out of scope, but COMGR still needed to access the binary data later, causing a use-after-free. - The fix ensures that the HIPOCProgramImpl object owns the binary data for its entire lifetime - Both constructors now consistently store the binary data in the `binary` member variable (std::vector) - The uint8_t constructor converts the data to char format using iterator range construction - This prevents the use-after-free that occurred when COMGR tried to access freed memory ## Test Plan Test output before change: ``` HSA_XNACK=1 ASAN_OPTIONS=symbolize=1 ./build/ml-libs/MIOpen/build/bin/miopen_gtest --gtest_filter="*CPU_Dump_NAN_FP32*" PRNG seed: 12345678 Note: Google Test filter = *CPU_Dump_NAN_FP32* [==========] Running 1 test from 1 test suite. [----------] Global test environment set-up. [----------] 1 test from CPU_Dump_NAN_FP32 [ RUN ] CPU_Dump_NAN_FP32.testDump ================================================================= ==3639==ERROR: AddressSanitizer: heap-use-after-free on address 0x7e0f08c50200 at pc 0x7f5f8d7a6554 bp 0x7ffcb7c4a730 sp 0x7ffcb7c49ee8 READ of size 26088 at 0x7e0f08c50200 thread T0 #0 0x7f5f8d7a6553 in memcpy /data/nhanna/repos/TheRock/compiler/amd-llvm/compiler-rt/lib/asan/../sanitizer_common/sanitizer_common_interceptors_memintrinsics.inc:117:5 ROCm#1 0x7f5f23d61d78 in COMGR::setCStr(char*&, llvm::StringRef, unsigned long*) /data/nhanna/repos/TheRock/compiler/amd-llvm/amd/comgr/src/comgr.cpp:216:9 ROCm#2 0x7f5f23d61d78 in COMGR::DataObject::setData(llvm::StringRef) /data/nhanna/repos/TheRock/compiler/amd-llvm/amd/comgr/src/comgr.cpp:334:17 ROCm#3 0x7f5f23d61d78 in amd_comgr_set_data /data/nhanna/repos/TheRock/compiler/amd-llvm/amd/comgr/src/comgr.cpp:606:24 ROCm#4 0x7f5f221dc1d3 in amd::Comgr::set_data(amd_comgr_data_s, unsigned long, char const*) /data/nhanna/repos/TheRock/rocm-systems/projects/clr/rocclr/device/comgrctx.hpp:252:12 ROCm#5 0x7f5f221dc1d3 in amd::device::Program::getSymbolsFromCodeObj(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>>>>*, amd_comgr_symbol_type_s) const /data/nhanna/repos/TheRock/rocm-systems/projects/clr/rocclr/device/devprogram.cpp:2061:14 ROCm#6 0x7f5f219e6f7c in hip::DynCO::populateDynGlobalVars() /data/nhanna/repos/TheRock/rocm-systems/projects/clr/hipamd/src/hip_code_object.cpp:216:22 ROCm#7 0x7f5f219e8e6a in hip::DynCO::getDynFunc(ihipModuleSymbol_t**, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>) /data/nhanna/repos/TheRock/rocm-systems/projects/clr/hipamd/src/hip_code_object.cpp:125:22 ROCm#8 0x7f5f21f842ba in hip::PlatformState::GetDynFunc(ihipModuleSymbol_t**, ihipModule_t*, char const*) /data/nhanna/repos/TheRock/rocm-systems/projects/clr/hipamd/src/hip_platform.cpp:884:22 ROCm#9 0x7f5f21ec2d71 in hip::hipModuleGetFunction(ihipModuleSymbol_t**, ihipModule_t*, char const*) /data/nhanna/repos/TheRock/rocm-systems/projects/clr/hipamd/src/hip_module.cpp:89:47 ROCm#10 0x7f5f2212c588 in hipModuleGetFunction /data/nhanna/repos/TheRock/rocm-systems/projects/clr/hipamd/src/hip_table_interface.cpp:1926:10 ROCm#11 0x7f5f7478d806 in miopen::HIPOCKernel::HIPOCKernel(miopen::HIPOCProgram, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>, std::vector<unsigned long, std::allocator<unsigned long>>, std::vector<unsigned long, std::allocator<unsigned long>>) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/include/miopen/hipoc_kernel.hpp:225:25 ROCm#12 0x7f5f766febb7 in miopen::KernelCache::AddKernel(miopen::Handle const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>, unsigned long, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, miopen::HIPOCProgram*) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/kernel_cache.cpp:161:18 ROCm#13 0x7f5f76b6f0e4 in miopen::Handle::AddKernel(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, unsigned long, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&) const /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/hip/handlehip.cpp:450:34 ROCm#14 0x7f5f7411b52f in miopen::checkNumericsImpl(miopen::Handle const&, int, miopen::TensorDescriptor const&, void const*, bool) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/check_numerics.cpp:107:12 ROCm#15 0x55e87c72ebee in void testDumpWithNan<float>(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/test/gtest/dumpTensorTest.cpp:130:8 ROCm#16 0x55e87c72d4e8 in CPU_Dump_NAN_FP32_testDump_Test::TestBody() /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/test/gtest/dumpTensorTest.cpp:157:37 ROCm#17 0x55e87ef19d5e in void testing::internal::HandleSehExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2653:27 ROCm#18 0x55e87ef19d5e in void testing::internal::HandleExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2689:52 ROCm#19 0x55e87ef04cdd in testing::Test::Run() /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2728:50 ROCm#20 0x55e87ef04cdd in testing::Test::Run() /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2718:6 ROCm#21 0x55e87ef04e64 in testing::TestInfo::Run() /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2874:14 ROCm#22 0x55e87ef0500e in testing::TestSuite::Run() /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:3052:33 ROCm#23 0x55e87ef0500e in testing::TestSuite::Run() /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:3006:6 ROCm#24 0x55e87ef0d20b in testing::internal::UnitTestImpl::RunAllTests() /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:6004:47 ROCm#25 0x55e87ef1a1de in bool testing::internal::HandleSehExceptionsInMethodIfSupported<testing::internal::UnitTestImpl, bool>(testing::internal::UnitTestImpl*, bool (testing::internal::UnitTestImpl::*)(), char const*) /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2653:27 ROCm#26 0x55e87ef1a1de in bool testing::internal::HandleExceptionsInMethodIfSupported<testing::internal::UnitTestImpl, bool>(testing::internal::UnitTestImpl*, bool (testing::internal::UnitTestImpl::*)(), char const*) /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2689:52 ROCm#27 0x55e87ef051b5 in testing::UnitTest::Run() /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:5583:55 ROCm#28 0x55e87eee3f9b in RUN_ALL_TESTS() /data/nhanna/repos/TheRock/build/third-party/googletest/dist/include/gtest/gtest.h:2334:73 ROCm#29 0x55e87eee3f9b in main /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/test/gtest/main_hip.cpp:34:12 ROCm#30 0x7f5f20a587e4 in __libc_start_main (/lib64/libc.so.6+0x3a7e4) (BuildId: 889235a2805b8308b2d0274921bbe1890e9a1986) ROCm#31 0x55e87b0bcf2d in _start (/data/nhanna/repos/TheRock/build/ml-libs/MIOpen/build/bin/miopen_gtest+0x126bf2d) 0x7e0f08c50200 is located 0 bytes inside of 26088-byte region [0x7e0f08c50200,0x7e0f08c567e8) freed by thread T0 here: #0 0x7f5f8d7b8ba2 in operator delete(void*, unsigned long) /data/nhanna/repos/TheRock/compiler/amd-llvm/compiler-rt/lib/asan/asan_new_delete.cpp:190:3 ROCm#1 0x7f5f76b7317d in std::__new_allocator<char>::deallocate(char*, unsigned long) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/new_allocator.h:172:2 ROCm#2 0x7f5f76b7317d in std::allocator<char>::deallocate(char*, unsigned long) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/allocator.h:210:25 ROCm#3 0x7f5f76b7317d in std::allocator_traits<std::allocator<char>>::deallocate(std::allocator<char>&, char*, unsigned long) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/alloc_traits.h:517:13 ROCm#4 0x7f5f76b7317d in std::_Vector_base<char, std::allocator<char>>::_M_deallocate(char*, unsigned long) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/stl_vector.h:390:4 ROCm#5 0x7f5f76b7317d in std::_Vector_base<char, std::allocator<char>>::~_Vector_base() /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/stl_vector.h:369:2 ROCm#6 0x7f5f76b7317d in std::vector<char, std::allocator<char>>::~vector() /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/stl_vector.h:738:7 ROCm#7 0x7f5f76b7317d in miopen::Handle::LoadProgram(std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, bool) const /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/hip/handlehip.cpp:633:5 ROCm#8 0x7f5f766fda82 in miopen::KernelCache::AddKernel(miopen::Handle const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>, unsigned long, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, miopen::HIPOCProgram*)::'lambda'()::operator()() const /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/kernel_cache.cpp:143:30 ROCm#9 0x7f5f766fda82 in miopen::KernelCache::AddKernel(miopen::Handle const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>, unsigned long, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, miopen::HIPOCProgram*) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/kernel_cache.cpp:124:26 ROCm#10 0x7f5f76b6f0e4 in miopen::Handle::AddKernel(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, unsigned long, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&) const /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/hip/handlehip.cpp:450:34 ROCm#11 0x7f5f7411b52f in miopen::checkNumericsImpl(miopen::Handle const&, int, miopen::TensorDescriptor const&, void const*, bool) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/check_numerics.cpp:107:12 ROCm#12 0x55e87c72ebee in void testDumpWithNan<float>(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/test/gtest/dumpTensorTest.cpp:130:8 ROCm#13 0x55e87c72d4e8 in CPU_Dump_NAN_FP32_testDump_Test::TestBody() /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/test/gtest/dumpTensorTest.cpp:157:37 ROCm#14 0x55e87ef19d5e in void testing::internal::HandleSehExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2653:27 ROCm#15 0x55e87ef19d5e in void testing::internal::HandleExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2689:52 previously allocated by thread T0 here: #0 0x7f5f8d7b7f9d in operator new(unsigned long) /data/nhanna/repos/TheRock/compiler/amd-llvm/compiler-rt/lib/asan/asan_new_delete.cpp:109:35 ROCm#1 0x7f5f76b720a5 in std::__new_allocator<char>::allocate(unsigned long, void const*) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/new_allocator.h:151:27 ROCm#2 0x7f5f76b720a5 in std::allocator<char>::allocate(unsigned long) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/allocator.h:198:32 ROCm#3 0x7f5f76b720a5 in std::allocator_traits<std::allocator<char>>::allocate(std::allocator<char>&, unsigned long) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/alloc_traits.h:482:20 ROCm#4 0x7f5f76b720a5 in std::_Vector_base<char, std::allocator<char>>::_M_allocate(unsigned long) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/stl_vector.h:381:20 ROCm#5 0x7f5f76b720a5 in std::_Vector_base<char, std::allocator<char>>::_M_create_storage(unsigned long) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/stl_vector.h:398:33 ROCm#6 0x7f5f76b720a5 in std::_Vector_base<char, std::allocator<char>>::_Vector_base(unsigned long, std::allocator<char> const&) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/stl_vector.h:335:9 ROCm#7 0x7f5f76b720a5 in std::vector<char, std::allocator<char>>::vector(std::vector<char, std::allocator<char>> const&) /opt/rh/gcc-toolset-13/root/usr/lib/gcc/x86_64-redhat-linux/13/../../../../include/c++/13/bits/stl_vector.h:602:9 ROCm#8 0x7f5f76b720a5 in miopen::Handle::LoadProgram(std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, bool) const /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/hip/handlehip.cpp:623:27 ROCm#9 0x7f5f766fda82 in miopen::KernelCache::AddKernel(miopen::Handle const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>, unsigned long, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, miopen::HIPOCProgram*)::'lambda'()::operator()() const /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/kernel_cache.cpp:143:30 ROCm#10 0x7f5f766fda82 in miopen::KernelCache::AddKernel(miopen::Handle const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>>, unsigned long, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, miopen::HIPOCProgram*) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/kernel_cache.cpp:124:26 ROCm#11 0x7f5f76b6f0e4 in miopen::Handle::AddKernel(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::filesystem::__cxx11::path const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::vector<unsigned long, std::allocator<unsigned long>> const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&, unsigned long, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&) const /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/hip/handlehip.cpp:450:34 ROCm#12 0x7f5f7411b52f in miopen::checkNumericsImpl(miopen::Handle const&, int, miopen::TensorDescriptor const&, void const*, bool) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/src/check_numerics.cpp:107:12 ROCm#13 0x55e87c72ebee in void testDumpWithNan<float>(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char>> const&) /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/test/gtest/dumpTensorTest.cpp:130:8 ROCm#14 0x55e87c72d4e8 in CPU_Dump_NAN_FP32_testDump_Test::TestBody() /data/nhanna/repos/TheRock/rocm-libraries/projects/miopen/test/gtest/dumpTensorTest.cpp:157:37 ROCm#15 0x55e87ef19d5e in void testing::internal::HandleSehExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2653:27 ROCm#16 0x55e87ef19d5e in void testing::internal::HandleExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) /data/nhanna/repos/TheRock/build/third-party/googletest/source/googletest/src/gtest.cc:2689:52 SUMMARY: AddressSanitizer: heap-use-after-free /data/nhanna/repos/TheRock/compiler/amd-llvm/amd/comgr/src/comgr.cpp:216:9 in COMGR::setCStr(char*&, llvm::StringRef, unsigned long*) Shadow bytes around the buggy address: 0x7e0f08c4ff80: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa 0x7e0f08c50000: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa 0x7e0f08c50080: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa 0x7e0f08c50100: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa 0x7e0f08c50180: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa =>0x7e0f08c50200:[fd]fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd 0x7e0f08c50280: fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd 0x7e0f08c50300: fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd 0x7e0f08c50380: fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd 0x7e0f08c50400: fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd 0x7e0f08c50480: fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd Shadow byte legend (one shadow byte represents 8 application bytes): Addressable: 00 Partially addressable: 01 02 03 04 05 06 07 Heap left redzone: fa Freed heap region: fd Stack left redzone: f1 Stack mid redzone: f2 Stack right redzone: f3 Stack after return: f5 Stack use after scope: f8 Global redzone: f9 Global init order: f6 Poisoned by user: f7 Container overflow: fc Array cookie: ac Intra object redzone: bb ASan internal: fe Left alloca redzone: ca Right alloca redzone: cb ==3639==ABORTING ``` ## Test Result Test output after change: ``` HSA_XNACK=1 ASAN_OPTIONS=symbolize=1 ./build/ml-libs/MIOpen/build/bin/miopen_gtest --gtest_filter="*CPU_Dump_NAN_FP32*" PRNG seed: 12345678 Note: Google Test filter = *CPU_Dump_NAN_FP32* [==========] Running 1 test from 1 test suite. [----------] Global test environment set-up. [----------] 1 test from CPU_Dump_NAN_FP32 [ RUN ] CPU_Dump_NAN_FP32.testDump [ OK ] CPU_Dump_NAN_FP32.testDump (51 ms) [----------] 1 test from CPU_Dump_NAN_FP32 (51 ms total) [----------] Global test environment tear-down [==========] 1 test from 1 test suite ran. (52 ms total) [ PASSED ] 1 test. ``` ## Cline Analysis ### Test Coverage Analysis: __1. LoadProgram Code Path (std::vector constructor):__ - __Primary Test__: `rocm-libraries/projects/miopen/test/gtest/db_sync.cpp` - __Function__: `BuildKernel()` calls `handle.LoadProgram(program_file, program_args, "")` - __Coverage__: This test extensively exercises the LoadProgram → LoadBinary → HIPOCProgramImpl constructor path - __Scope__: Tests multiple GPU architectures (gfx908, gfx90a, gfx942, gfx1030) with different CU counts - __Frequency__: Runs on thousands of kernel configurations in the database sync tests __2. Solution Binary Serialization (std::vector usage):__ - __Primary Test__: `rocm-libraries/projects/miopen/test/gtest/find_2_conv.cpp` - __Function__: `miopenSaveSolution()` and `miopenLoadSolution()` with `std::vector<char> solution_binary` - __Coverage__: Tests the save/load cycle of solution binaries - __Scope__: Tests all convolution directions (Forward, BackwardData, BackwardWeights) __3. Additional Coverage:__ - __Cache Tests__: `rocm-libraries/projects/miopen/test/gtest/cache.cpp` tests compression/decompression with `std::vector<char>` - __Dropout Tests__: Uses `std::vector<unsigned char>` for reserve space (related pattern) __Test Quality Assessment:__ ✅ __Both constructors are well-tested__: - The `std::vector<char>` constructor is heavily exercised through database sync tests - The `std::vector<uint8_t>` constructor would be tested through any code paths that use uint8_t binary data ✅ __Real-world scenarios covered__: - Database synchronization (production kernel loading) - Solution serialization (runtime binary handling) - Multi-threaded execution (db_sync uses up to 32 threads) ✅ __Comprehensive architecture coverage__: - Tests run on multiple GPU architectures - Different compute unit configurations tested __Confidence Level__: Very High ### Performance Analysis: Regarding the performance impact of this fix, it's actually quite minimal and represents good engineering practice: __Memory Impact:__ - __Additional Memory Usage__: Each HIPOCProgramImpl object now stores a copy of the binary data in its `binary` member variable - __Typical Size__: GPU code objects are usually relatively small (typically a few KB to a few MB depending on kernel complexity) - __Lifetime__: The memory is only held for the lifetime of the HIPOCProgram object, which is typically short-lived during kernel loading __Performance Characteristics:__ - __One-time Copy Cost__: There's a single memory copy operation during construction (std::vector copy or iterator range construction) - __No Runtime Overhead__: Once constructed, there's no additional performance cost during kernel execution - __Memory Safety Benefit__: Eliminates potential crashes and undefined behavior, which far outweighs the small memory cost __Context in MIOpen:__ - This occurs during the kernel loading phase, not during actual ML inference/training - Kernel loading is already an expensive operation involving compilation, module creation, etc. - The additional memory copy is negligible compared to the overall kernel loading time __Trade-off Analysis:__ - __Cost__: Small increase in memory usage during kernel loading - __Benefit__: Eliminates memory safety bugs that could cause crashes or data corruption - __Net Result__: Significantly positive - reliability and correctness are much more valuable than the minimal memory overhead In practice, this fix follows the RAII (Resource Acquisition Is Initialization) principle and ensures proper ownership semantics, which is standard best practice in modern C++. The performance impact should be unnoticeable in real-world usage.

…7008) ## Motivation Several memory leaks were detected in MIOpen gtests using ASan. Some of the tests were blacklisted and others were not. This change looks to fix all of the low hanging fruit, which are the majority of the leaks found. This includes all of the critical leaks (>100MB) that were reported. Some other leaks were identified as needing a larger refactor to resolve. After fixing the w_supertensor.cpp leak, the supertensor tests hit virtual memory area limit errors. These changes are to enable running them with ASan, but only with a subset of the tests in order to not hit limits. The test coverage being lost is fairly negligible and the full tests are still run without ASan enabled. ## Technical Details Here is a summary of the files looked at and the changes made: | File | Status | Notes | |------|--------|-------| | `test/gtest/gpu_mha_forward.cpp` | Fixed | Added `miopenDestroySolution()` loop after using solutions. Fixes MHA forward solution descriptor leaks (report #11, #13). | | `test/gtest/gpu_mha_backward.cpp` | Fixed | Added `miopenDestroySolution()` loop after using solutions. Fixes MHA backward solution descriptor leaks (report #9, #12). | | `test/gtest/mha_find20.cpp` | Fixed | Added `miopenDestroySolution()` loop in both `MhaForward` and `MhaBackward` tests. Fixes MHA Find2.0 solution leak (report #16). | | `test/gtest/gtest_desc_guard.hpp` | Fixed | New shared header introduced by the refactor. Provides a single `DescGuard<DescType, CreateFn, DestroyFn>` template (with `TensorDescGuard`, `ConvDescGuard`, `DropoutDescGuard`, `RNNDescGuard` aliases), a `HandleGuard` RAII wrapper for `miopenHandle_t`, and the `DestroyInternalRnnDropoutDesc(rnnDesc)` helper used by every RNN/LSTM/GRU test to free the internal `DropoutDescriptor` that `miopenCreateRNNDescriptor` allocates and `miopenSetRNNDescriptor*` then leaks. Replaces the per-file ad-hoc guard structs from the initial implementation. | | `test/gtest/w_supertensor.cpp` | Fixed | Switched raw descriptors to the shared `RNNDescGuard` / `TensorDescGuard` from `gtest_desc_guard.hpp`. Added a class-local `DestroyDropoutDesc()` (called from `TearDown` and before `miopenSetRNNDescriptor`) to prevent the `miopenSetRNNDescriptor` overwrite leak. Reduced test parameter space under ASan to avoid OOM. Removed unused `seqLen` parameter and dead `param_dev_out`/`bias_dev_out` allocations. | | `test/gtest/lstm.hpp` | Fixed | Switched `rnnDesc` → `RNNDescGuard`, `DropoutDesc` → `DropoutDescGuard`, and `mio_handle` → `HandleGuard` (which now owns the `miopenDestroy` call), all from the shared `gtest_desc_guard.hpp`. Hoisted `dropout_state_buf` so it can be `hipFree`d at the end of the dropout path. Added `DestroyInternalRnnDropoutDesc(rnnDesc)` before `miopenSetRNNDescriptor*` and (in the non-dropout path only) at end of `Run`, which frees the internal `DropoutDescriptor` that the Set call would otherwise leak. Fixes ~615 MB dropout leaks and ~16.5 KB non-dropout descriptor leaks (report #3, #4, #6, #17-#22). | | `test/gtest/gru_test.cpp` | Fixed | In the GRU test class: switched `rnnDesc` → `RNNDescGuard`, `DropoutDesc` → `DropoutDescGuard`, `mio_handle` → `HandleGuard`, and added `DestroyInternalRnnDropoutDesc(rnnDesc)` before `miopenSetRNNDescriptor*` and (non-dropout path only) at end of `Run` plus `hipFree(dropout_state_buf)` for the dropout path. In the in-file `GRUFwdCPUVerify` / `GRUBwdDataCPUVerify` helpers: converted the raw `dropout_inputTensor` / `dropout_outputTensor` declarations to `TensorDescGuard` (mirroring the `cpu_rnn.hpp` change for the LSTM/RNN helpers). | | `test/gtest/softmax_find20.cpp` | Fixed | Changed `Finalize()` to take the `std::vector<miopenSolution_t>&` and destroy each solution via `miopenDestroySolution()` before destroying the problem. Updated all 6 `TEST(...)` callers to pass the solutions vector. Fixes Find2.0 softmax solution/kernel leaks (report #25-#27). | | `test/gtest/rnn_seq_api.hpp` | Fixed | Hoisted `dropout_state_buf` so the dropout path can `hipFree` it at the end. Added `DestroyInternalRnnDropoutDesc(&rnnDesc)` before `miopenSetRNNDescriptor*` and (non-dropout path only) at end of the run to free the internal `DropoutDescriptor` allocations leaked by `RNNDescriptor` copy-assignment. Same pattern as the LSTM/GRU fixes (report #29-#30). | | `test/cpu_rnn.hpp` | Fixed | Converted 6 raw `miopenTensorDescriptor_t dropout_input/outputTensor` declarations across the LSTM/RNN CPU verification helpers (`LSTMFwdCPUVerify`, `LSTMBwdDataCPUVerify`, `RNNFwdTrainCPUVerify`, `RNNBwdDataCPUVerify`, `GRUFwdCPUVerify`, `GRUBwdDataCPUVerify`) to the shared `TensorDescGuard` from `gtest_desc_guard.hpp`. Removed the redundant `miopenCreateTensorDescriptor` calls and updated 12 `miopen::deref(...)` sites to `.get()`. (Note: the GRU helpers in this header are stale duplicates; the live ones are inside `test/gtest/gru_test.cpp` and were updated there too.) Fixes the LSTM/GRU CPU-verify tensor descriptor leaks (report #1, #2, #14, #15, #32, #33). | | `test/gtest/rnn_vanilla_common.hpp` | Fixed | Added `DestroyInternalRnnDropoutDesc(rnnDesc)` calls before `miopenSetRNNDescriptor*` and (non-dropout path only) at end of the run. The `RNNDescGuard` / `DropoutDescGuard` usage was already in place from an earlier commit and now resolves to the shared definitions in `gtest_desc_guard.hpp`. Same pattern as LSTM/GRU/rnn_seq_api fixes (report #14, #15, #32, #33). | | `test/gtest/graphapi_gtest_common.hpp` | Skipped | File no longer exists in the codebase. The GraphApi test infrastructure has been removed; only the leak report (a stale snapshot) still references it. No fix possible against the current source tree. | | `test/gtest/graphapi_execution_plan.cpp` | Skipped | File no longer exists in the codebase (GraphApi removed). The leak it represented was largely an external hipblaslt bug anyway; the remaining test-side portion is not fixable against the current source. | | `test/gtest/na_train.cpp` / `na_inference.cpp` / `na_*_find2.cpp` | Skipped | Leaks come from the internal MIOpen solver/kernel cache living on the global singleton handle, which is never destroyed. Not easily fixable without redesigning global handle lifecycle. Should be suppressed in the ASan suppression file. | | hipblaslt / rocblaslt (external) | Skipped | `SolutionCache::addKernel` and `preloadCustomKernels` leak via `_rocblaslt_handle` constructor. Called from `miopen::Handle::CreateHipblasLtHandle`. This is an upstream bug in hipblaslt/rocblaslt, not fixable in MIOpen. Affects suites that initialize a handle (report Category 2, Category 7). | | CLR / HIP runtime (external) | Skipped | `amd::Context` and `amd::roc::Device` global initialization leaks from `rocclr/platform/context.cpp`. HIP runtime internals, not fixable in MIOpen. | | `src/hipoc/hipoc_program.cpp` | Skipped | `HIPOCProgramImpl` objects leak during kernel compilation/caching (line ~178). This is an internal MIOpen kernel cache lifecycle issue that requires deeper architectural changes to fix. Contributes small amounts to MHA and Softmax Find2.0 leaks. | | `test/gtest/conv_api.cpp` | Fixed | Already clean against current source — `miopenDestroyConvolutionDescriptor(conv_desc)` call exists (line 24) inside the test loop. ASAN run reports no leaks. The leak report was based on a stale snapshot. The hipblaslt handle init portion is tracked under the external-skipped row. | | `test/gtest/log_test.cpp` (CPU_LOG_TEST_FUSION / CPU_LOG_TEST_NEG) | Fixed | Already clean against current source — `Tensor`, `Conv`, `CreateCBAFusionPlan`, `CreateBNormFusionPlan` all have proper destructors that call the corresponding `miopenDestroy*` APIs in `log.cpp`. ASAN run on `CPU_LOG_TEST_*` (11 tests across log_test.cpp + log_test_neg.cpp) reports no leaks. The hipblaslt handle init portion is tracked separately under the external-skipped row. | | `test/gtest/fusion_test.cpp` (CPU_FusionCreateOpConvForward) | Fixed | File renamed from `fusion.cpp` to `fusion_test.cpp`. Already clean against current source — uses `TensorDescGuard`/`ConvDescGuard` for tensor/conv descriptors and calls `miopenDestroyFusionPlan(fusionPlanDesc)` on the fusion plan (line 195). ASAN run on `CPU_FusionCreateOpConvForward_FP32.*` reports no leaks. | | `test/gtest/deterministic_conv_api.cpp` | Fixed | Already clean against current source — uses `ConvDescGuard` (line 66) for the conv descriptor. ASAN run on `*CPU_DeterministicConvApi*` reports no leaks. | | `test/gtest/fusion_aux.cpp` (GPU_FusionAux) | Fixed | Already clean against current source — uses `ConvDescGuard` plus stack-allocated internal C++ objects (`miopen::TensorDescriptor`, `miopen::FusionPlanDescriptor`) which have proper destructors. The `convoOp` handle is owned by the fusion plan. ASAN run on `*GPU_FusionAux*` reports no leaks. | | `test/gtest/backend_api.cpp` (CPU_BackendApi) | Skipped | File no longer exists in the codebase. The backend API test infrastructure (part of the removed GraphApi suite) was removed; no fix possible against the current source tree. | ### High-level notes New shared infrastructure (test/gtest/gtest_desc_guard.hpp) - DescGuard<DescType, CreateFn, DestroyFn> — a single RAII template parameterized on the descriptor type and its create/destroy entry points. Aliases provide TensorDescGuard, ConvDescGuard, DropoutDescGuard, and RNNDescGuard, replacing the four near-identical guard structs that were copy-pasted across test files in the initial implementation. - HandleGuard — separate RAII wrapper for miopenHandle_t (couldn't reuse the template because miopenCreateWithStream takes an extra hipStream_t argument). Supports lazy create(stream) so callers that only need a handle in the dropout branch can default-construct one and populate it conditionally. - DestroyInternalRnnDropoutDesc(rnnDesc) — frees the internal DropoutDescriptor that miopenCreateRNNDescriptor allocates and that miopenSetRNNDescriptor* then orphans. Replaces the equivalent inline blocks that LSTM/GRU/RNN tests were each carrying. The header documents the two call-sites: before each Set* (always safe) and at end-of-run only on the non-dropout path (the dropout path aliases the user-owned descriptor, so freeing would double-free). Recurring patterns enabled by the refactor - The "leak from Set* overwriting the default-constructed internal dropout descriptor" fix collapsed from per-file code to a one-line helper call, applied uniformly across lstm.hpp, gru_test.cpp, rnn_seq_api.hpp, and rnn_vanilla_common.hpp. - mio_handle ownership in LSTM/GRU is now expressed via HandleGuard rather than a manual miopenDestroy at the end of the dropout branch — eliminates a class of forgotten-cleanup bugs. - dropout_state_buf is consistently hoisted out of the dropout if block so an end-of-run hipFree can release it; deletion of the buffer pairs visibly with its allocation. Notable non-RNN change - softmax_find20.cpp was the only Find2.0 leak fix in this commit: Finalize() now takes the solutions vector and calls miopenDestroySolution() for each before destroying the problem. Same shape applied to all 6 tests in the file. ## Test Plan Run the tests beforehand to observe the ASan leak errors and then again afterward to verify the fixes have resolved the problem. ## Test Result List from ROCM-21512: | # | Test Name | Status | Leak Status | |---:|------------------------------------------------------|------------------------------------------------|-----------------| | 1 | Smoke/GPU_RNNVanillaDropout_FP32 | PASSED (4 tests) | No leaks | | 2 | Smoke/GPU_RNNVanillaDropout_FP16 | PASSED (4 tests) | No leaks | | 3 | Full/GPU_LSTM_dropout_FP32 | PASSED (4 tests) | No leaks | | 4 | Full/GPU_LSTM_dropout_FP16 | PASSED (4 tests) | No leaks | | 5 | CPU_GraphApiExecutionPlanBuilder_NONE | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 6 | Full/GPU_LSTM_dropout_FP64 | REMOVED (PR #5750, 2026-03-26) | n/a (deleted) | | 7 | Unit/CPU_GraphApiPointwise_NONE | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 8 | Full/GPU_LstmMSRnn_FP32 | PASSED (1152 tests) | No leaks | | 9 | Smoke/GPU_Bwd_Mha_FP32 | PASSED (12 tests) | No leaks | | 10 | Full/GPU_LstmMSRnn_FP16 | PASSED (864 tests) | No leaks | | 11 | Smoke/GPU_Fwd_Mha_FP32 | PASSED (15 tests) | No leaks | | 12 | Full/GPU_Bwd_Mha_FP32 | PASSED (6 tests) | No leaks | | 13 | Full/GPU_Fwd_Mha_FP32 | PASSED (7 tests) | No leaks | | 14 | Full/GPU_RNNVanilla_FP32 | PASSED (96 tests) | No leaks | | 15 | Full/GPU_RNNVanilla_FP16 | PASSED (96 tests) | No leaks | | 16 | GPU_TestMhaFind20_FP32 | PASSED (2 tests) | No leaks | | 17 | Full/GPU_LSTM_FP32 | PASSED (32 tests) | No leaks | | 18 | Full/GPU_LSTM_FP16 | PASSED (32 tests) | No leaks | | 19 | Full/GPU_LSTM_extra_FP32 | PASSED (30 tests) | No leaks | | 20 | Full/GPU_LSTM_extra_FP16 | PASSED (30 tests) | No leaks | | 21 | Full/GPU_DeepBench_LSTM_FP16 | PASSED (22 tests) | No leaks | | 22 | Full/GPU_DeepBench_LSTM_FP32 | PASSED (22 tests) | No leaks | | 23 | CPU_LOG_TEST_FUSION_NONE | PASSED (2 tests) | No leaks | | 24 | CPU_LOG_TEST_NEG_NONE | PASSED (4 tests) | No leaks | | 25 | GPU_SoftmaxFind20_BFP16 | PASSED (2 tests) | No leaks | | 26 | GPU_SoftmaxFind20_FP16 | PASSED (2 tests) | No leaks | | 27 | GPU_SoftmaxFind20_FP32 | PASSED (2 tests) | No leaks | | 28 | CPU_ConvApi_NONE | PASSED (1 test) | No leaks | | 29 | Full/GPU_RNNSeqApi_FP16 | PASSED (16 tests) | No leaks | | 30 | Full/GPU_RNNSeqApi_FP32 | PASSED (16 tests) | No leaks | | 31 | UnitVAN/CPU_GraphApiRng_NONE | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 32 | Smoke/GPU_RNNVanilla_FP16 | PASSED (4 tests) | No leaks | | 33 | Smoke/GPU_RNNVanilla_FP32 | PASSED (4 tests) | No leaks | | 34 | CPU_FusionCreateOpConvForward_FP32 | PASSED (1 test) | No leaks | | 35 | CPU_GraphApiOperationReduction_NONE | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 36 | Unit2IV1/CPU_GraphApiOperationPointwiseBwd_NONE | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 37 | Unit2IV1/CPU_GraphApiOperationPointwiseTwoInput | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 38 | Unit2IV2/CPU_GraphApiOperationPointwiseBwd_NONE | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 39 | Unit2IV2/CPU_GraphApiOperationPointwiseTwoInput | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 40 | UnitVAB/CPU_GraphApiRng_NONE | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 41 | Smoke/GPU_RNNVanillaDropout_FP16 (duplicate of #2) | (see #2) | (see #2) | | 42 | CPU_GraphApiOperationGraphDescriptor_NONE | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 43 | UnitVA/CPU_GraphApiVariantPack_NONE | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 44 | UnitVAU/CPU_GraphApiRng_NONE | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 45 | CPU_GraphApiOperationReshape_NONE | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 46 | Smoke/CPU_DeterministicConvApi_NONE | PASSED (1 test) | No leaks | | 47 | Smoke/GPU_FusionAux_FP32 | PASSED (1 test) | No leaks | | 48 | CPU_GraphApiEngineHeur_NONE | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 49 | Unit/CPU_GraphApiReduction_NONE | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 50 | CPU_GraphApiEngineCfg_NONE | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 51 | Unit/CPU_GraphApiMatMul_NONE | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 52 | CPU_BackendApi_NONE | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 53 | UnitIV/CPU_GraphApiOperationPointwiseOneInput | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 54 | Unit3IV/CPU_GraphApiOperationPointwiseThreeInput | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 55 | UnitVA/CPU_GraphApiOperationMatmul_NONE | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | 56 | UnitVA/CPU_GraphApiOperationRng_NONE | REMOVED (PR #5603, 2026-03-26) | n/a (deleted) | | Outcome | Count | |------------------------------------|------:| | Passed, no leaks | 39 | | Passed, leaks detected | 0 | | Failed | 0 | | Crashed / timed out | 0 | | Removed — GraphAPI purge (#5603) | 15 | | Removed — FP64 LSTM purge (#5750) | 1 | | Duplicate (not re-run) | 1 | | **Total rows** | **56** | ## Risk Assessment Low --------- Co-authored-by: JonathanLichtnerAMD <195780826+JonathanLichtnerAMD@users.noreply.github.com>

…rom full ctest Two fixes from observing build #27: 1. The per-arch Smart Build stages each publish an identical smoke_result.xml, so Jenkins shows 5 indistinguishable "selection-resolvable" rows under "full". Thread ARCH_NAME into the JUnit suite/classname (smart-build.selection.full.gfx942) via a new validate --label, so per-arch rows are distinct. 2. Full-mode smart_test ran the whole ctest registry including rocm_ck_unit, which is registered with ctest but built by check-rocm-ck (not ninja tests examples), so ctest reported it "Not Run" and failed the build. Exclude the separate flag-gated suites (ROCM_CK_*, BUILDER_SMOKE) via ctest -LE in full mode (D13), overridable via CTEST_FULL_EXCLUDE_LABELS. Generated-by: Claude Code (claude-opus-4-8)

[GitHub Actions] Changing default behaviour for PR imports

004bece

jayhawk-commits merged commit 7138fc8 into develop May 13, 2025

jayhawk-commits deleted the users/jayhawk-commits/import-defaults branch May 13, 2025 22:49

jayhawk-commits self-assigned this May 17, 2025

jayhawk-commits added github actions migration Tasks or issues tied to migration to this monorepo labels May 17, 2025

assistant-librarian Bot mentioned this pull request Jul 14, 2025

rocsparse-test fails to run on Fedora #627

Closed

assistant-librarian Bot mentioned this pull request Jul 23, 2025

hipblas-test double free error on Fedora. #807

Closed

evetsso pushed a commit to evetsso/rocm-libraries that referenced this pull request Dec 31, 2025

Merge pull request ROCm#27 from AMD-ROCm-Internal/lima1002/gfx1250-wmma

3908631

Enable wmma-related tests on gfx1250

NolanHannaAMD mentioned this pull request May 1, 2026

[MIOpen] [ROCM-21512] [ALMIOPEN-1814] Fix MIOpen ASan memory leaks #7008

Merged

Navigation Menu

Uh oh!

[GitHub Actions] Changing default behaviour for PR imports - #27

[GitHub Actions] Changing default behaviour for PR imports#27
jayhawk-commits merged 1 commit into
developfrom
users/jayhawk-commits/import-defaults

jayhawk-commits commented May 13, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Uh oh!

Conversation

jayhawk-commits commented May 13, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant