Testing auto-label job post folder reorg by jayhawk-commits · Pull Request #7 · ROCm/rocm-libraries

jayhawk-commits · 2025-04-30T22:46:47Z

No description provided.

…into testAutoLabel

* support 16x16x128 f8 wmma * wip: use 128-bit ds load for 8-bit wmma * Added test yaml for fp8 wmma * Force to use ds_load_b128 for gfx1250 * Removed duplicated f8f6f4 wmma capability check * Removed un-used testing yaml * Reduced number of tests of gfx1250 fp8 wmma

* Adding initial data schema for graphs * Add Node definition * Add frontend API for attributes * Replace frontend tests with frontend api tests * Remove vibe comments * Add dependency on SDK target to ensure generation happens * Restructure directories * Formatting

* Adding initial data schema for graphs * Add Node definition * Add frontend API for attributes * Replace frontend tests with frontend api tests * Remove vibe comments * Add dependency on SDK target to ensure generation happens * Restructure directories * Formatting [ROCm/hipDNN commit: d8d0f29]

## 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.

#3710) ## Motivation Optimizing the tensor filling functions started a discussion about optimizing tensor iteration in general: #3471 (comment) ## Technical Details After some deliberation, the approach taken here (using std::variant inside the iterator to represent the different types of indexing) reflects both the desire the improve iteration in the case of packed tensors while also maintaining the existing API. A fully templated approach would be more optimal but would require API changes to the ITensor class itself, whether making it templated or changing the definition of its iterator-related methods at the very least. ## Test Plan Ran ninja check inside the build folder of hipDNN. ## Test Result ``` [185/187] Running all tests via ctest Test project /therock/output/build/ml-libs/hipDNN/build Start 1: hipdnn_data_sdk_tests 1/7 Test #1: hipdnn_data_sdk_tests ............ Passed 1.30 sec Start 2: hipdnn_backend_tests 2/7 Test #2: hipdnn_backend_tests ............. Passed 1.29 sec Start 3: hipdnn_frontend_tests 3/7 Test #3: hipdnn_frontend_tests ............ Passed 0.03 sec Start 4: hipdnn_test_sdk_tests 4/7 Test #4: hipdnn_test_sdk_tests ............ Passed 4.32 sec Start 5: hipdnn_plugin_sdk_tests 5/7 Test #5: hipdnn_plugin_sdk_tests .......... Passed 0.03 sec Start 6: public_hipdnn_backend_tests 6/7 Test #6: public_hipdnn_backend_tests ...... Passed 0.33 sec Start 7: public_hipdnn_frontend_tests 7/7 Test #7: public_hipdnn_frontend_tests ..... Passed 0.26 sec 100% tests passed, 0 tests failed out of 7 Label Time Summary: integration_test = 0.59 sec*proc (2 tests) unit_test = 6.96 sec*proc (5 tests) Total Test time (real) = 7.56 sec ``` ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: BrianHarrisonAMD <169072757+BrianHarrisonAMD@users.noreply.github.com>

## Motivation  This PR addresses hipDNN issue #4951, which requests adding missing frontend integration test coverage for the Matmul operation. ## Technical Details  This PR includes three changes: ### 1. FrontendGraphFactory support for Matmu Added MATMUL to OperationType. Added switch‑case dispatch in FrontendGraphFactory::create(). Implemented createMatmulGraph() using: ```cpp graph.matmul(a, b, matmulAttrs); ``` with simple 2×3 and 3×4 matrix inputs for deterministic testing. ### 2. Added new integration test: IntegrationMatmul.cpp Following the structure of IntegrationConvForward.cpp, the test: - is parameterized with: - good plugin - execute‑fail plugin - no‑engines plugin - tests both auto‑assigned and manual UIDs - builds a small Matmul graph using float tensors - exercises the entire frontend execution pipeline: ``` validate() → build_operation_graph() → create_execution_plans() → check_support() → build_plans() → get_workspace_size() → execute() ``` - uses SKIP_IF_NO_DEVICES() for GPU‑dependent execution - creates variant packs using device memory from the test tensor bundle - verifies expected failures for execute‑fail and no‑engines plugins ### 3. CMake update Added IntegrationMatmul.cpp to tests/frontend/CMakeLists.txt under public_hipdnn_frontend_tests. ## Test Plan  All tests were built and executed inside the official TheRock docker environment. ## Test Result  ``` [1/2] Validating test names with --gtest_list_tests test collection Test Name Validation Report ============================================================ Total tests found: 2901 Valid test names: 2901 Invalid test names: 0 ``` ``` [1/2] Running all tests via ctest Test project /therock/output/build/ml-libs/hipDNN/build Start 1: hipdnn_data_sdk_tests 1/7 Test #1: hipdnn_data_sdk_tests ............ Passed 0.96 sec Start 2: hipdnn_backend_tests 2/7 Test #2: hipdnn_backend_tests ............. Passed 1.38 sec Start 3: hipdnn_frontend_tests 3/7 Test #3: hipdnn_frontend_tests ............ Passed 0.05 sec Start 4: hipdnn_test_sdk_tests 4/7 Test #4: hipdnn_test_sdk_tests ............ Passed 8.19 sec Start 5: hipdnn_plugin_sdk_tests 5/7 Test #5: hipdnn_plugin_sdk_tests .......... Passed 0.03 sec Start 6: public_hipdnn_backend_tests 6/7 Test #6: public_hipdnn_backend_tests ...... Passed 0.32 sec Start 7: public_hipdnn_frontend_tests 7/7 Test #7: public_hipdnn_frontend_tests ..... Passed 0.35 sec ``` ``` 100% tests passed, 0 tests failed out of 7 Label Time Summary: integration_test = 0.67 sec*proc (2 tests) unit_test = 10.61 sec*proc (5 tests) Total Test time (real) = 11.29 sec ``` ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Signed-off-by: jovanau <u.jovana2@gmail.com>

…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.

## Motivation  This PR addresses hipDNN issue ROCm#4951, which requests adding missing frontend integration test coverage for the Matmul operation. ## Technical Details  This PR includes three changes: ### 1. FrontendGraphFactory support for Matmu Added MATMUL to OperationType. Added switch‑case dispatch in FrontendGraphFactory::create(). Implemented createMatmulGraph() using: ```cpp graph.matmul(a, b, matmulAttrs); ``` with simple 2×3 and 3×4 matrix inputs for deterministic testing. ### 2. Added new integration test: IntegrationMatmul.cpp Following the structure of IntegrationConvForward.cpp, the test: - is parameterized with: - good plugin - execute‑fail plugin - no‑engines plugin - tests both auto‑assigned and manual UIDs - builds a small Matmul graph using float tensors - exercises the entire frontend execution pipeline: ``` validate() → build_operation_graph() → create_execution_plans() → check_support() → build_plans() → get_workspace_size() → execute() ``` - uses SKIP_IF_NO_DEVICES() for GPU‑dependent execution - creates variant packs using device memory from the test tensor bundle - verifies expected failures for execute‑fail and no‑engines plugins ### 3. CMake update Added IntegrationMatmul.cpp to tests/frontend/CMakeLists.txt under public_hipdnn_frontend_tests. ## Test Plan  All tests were built and executed inside the official TheRock docker environment. ## Test Result  ``` [1/2] Validating test names with --gtest_list_tests test collection Test Name Validation Report ============================================================ Total tests found: 2901 Valid test names: 2901 Invalid test names: 0 ``` ``` [1/2] Running all tests via ctest Test project /therock/output/build/ml-libs/hipDNN/build Start 1: hipdnn_data_sdk_tests 1/7 Test ROCm#1: hipdnn_data_sdk_tests ............ Passed 0.96 sec Start 2: hipdnn_backend_tests 2/7 Test ROCm#2: hipdnn_backend_tests ............. Passed 1.38 sec Start 3: hipdnn_frontend_tests 3/7 Test ROCm#3: hipdnn_frontend_tests ............ Passed 0.05 sec Start 4: hipdnn_test_sdk_tests 4/7 Test ROCm#4: hipdnn_test_sdk_tests ............ Passed 8.19 sec Start 5: hipdnn_plugin_sdk_tests 5/7 Test ROCm#5: hipdnn_plugin_sdk_tests .......... Passed 0.03 sec Start 6: public_hipdnn_backend_tests 6/7 Test ROCm#6: public_hipdnn_backend_tests ...... Passed 0.32 sec Start 7: public_hipdnn_frontend_tests 7/7 Test ROCm#7: public_hipdnn_frontend_tests ..... Passed 0.35 sec ``` ``` 100% tests passed, 0 tests failed out of 7 Label Time Summary: integration_test = 0.67 sec*proc (2 tests) unit_test = 10.61 sec*proc (5 tests) Total Test time (real) = 11.29 sec ``` ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Signed-off-by: jovanau <u.jovana2@gmail.com>

…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.

…nfigs (#7) * 64x64 * Fix MFMA emit code * Remove label * cleanup * cleanup * Cleanup * Update tests * New GR offset calculation (no swizzling yet) * Refacto * cleanup * Re-enable swizzling * Fix SPGR alloc * Update M0 * Tensile passing no swizzling * Fix swizzling * LDS padding * as is * Multiple bugfixes * Fix 128x64 * Refactor pre-swizzling change * Add wave specific rotation to swizzling * Fix gra Test * Fix LRA test * Fix roundtrip test * LdsNumBytes as int * Use float type for bpe * Cleanup * Cleanup * More cleanup * cleanup * Simplify _grSwizzleColIds * Remove debug label * Fix typo LDS size calculation

CK Tile's fmha_bwd.py get_dq_dk_dv_tiles() defines tuned dQ/dK/dV tile configs only for symmetric head dimensions (hdim_q == hdim_v). The prior table keyed on max(hdim_q, hdim_v) and synthesized configs for asymmetric pairs by extrapolation, which has no validated tile and could select an illegal GEMM4 warp distribution. Reject asymmetric configs instead. Addresses PR #7538 review findings: - Collapse generateTileConfig(hdim_q, hdim_v) to generateTileConfig(hdim); getTileConfig now throws on hdim_q != hdim_v as the single gate, so makeSpec's duplicate asymmetric reject-list is removed (#1, #4). - Keep the consteval derivation engine, symmetric-only (#3). - Add validateWaveTiles() so a future base-tile edit that yields an illegal MFMA/WMMA warp shape fails the build; uses the previously unused isValidWaveTile() as the single source of truth (#6). - Fix the "occupancy -1 = auto" doc contradiction and assert the table stores occupancy >= 1 (#2). - Remove the dead GEMM4 (4,1) warp branch; it is unreachable for the shipped gfx9 base tiles (#8). - Reword computeGemm4Warps as a divisibility rule, not a head-dim-size heuristic (#11). - Drop the now self-referential asymmetric tests from the consteval and compat suites; add two compile_fail/ negative TUs covering both asymmetric directions (#7, #9). Verified: positive smoke compiles; both negatives fail to compile with "requires hdim_q == hdim_v"; gtest green (consteval 21/21, dqdkdv 34/34); clang-format-18 applied. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

The tile-geometry and block_per_cu comments still claimed "d128 only" after d64 was wired. Clarify that the wired d64 and d128 rows both use bn0=128 (so the hardcoded block_n0 is correct for both, guarded by the device static_assert) and that production would derive these per hdim. Comment-only; no behavior change. Addresses review finding #7 on PR #7804. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

…, #12, stacked on #1, #2] Add double-precision GEMM support to the CPU GEMM driver: - Add TypeTraits<double> mapping to DataType::Double - Add --type f64 branch calling runGemm<double, double> - Type all auxiliary buffers (bias, scaleAlphaVec, scaleAB) as AccumulateT so the slow path's GetValue(alphaType, ...) reads the correct byte width when alphaType=Double - Set alphaType/betaType to accumDtypeEnum so TypeId dispatch matches TypedGemm_D_D_D for the slow reference solver - Register bias/scale tensors with accumDtypeEnum - Add 10 f64 slow-path tests: 4 transpose combos, Beta, Bias, AllFeatures, TN_AllFeatures, ScaleAB Scalar/Vector Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

The driver's FP4 init values previously covered only the small power-of-two slice {±0, ±0.5, ±1, ±2}, missing half of the E2M1-representable grid. Add the missing magnitudes (±1.5, ±3, ±4, ±6) so the MX-scale validation exercises products that span more of the FP4 range. All new values are exactly representable in FP4, so quantization is still a no-op and the existing tolerance holds.

Resolve the contradiction between the bundle-metadata section (which read as "verifier rejects bundles without the mandatory fields") and the tolerance framework (missing meta falls back to per-operation default, "no crash, no skip"). A forward bundle without .meta.json is valid; the mandatory fields are enforced only when a .meta.json is present. Backward bundles still require forward_source. Matches verifier check #7 (backward-only) and lines 522/579. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

Refresh the authoritative status docs now that the campaign is done: - PLAN-CODEGEN-WORKFLOW.md: rewrite the top "CURRENT STATE" header from "NEXT = P4 / 68.85%" to "P0-P6 COMPLETE, 80.70% (>=80% met)", with the full round trajectory, commit list, and optional follow-ups. - BASELINE-AND-PROGRESS.md: update the section-4 prose headline 68.85% -> 80.70% (final), close the "gap to 80%" subsection, and extend the receipt index with the P4/P5/P6 receipts + workflow scripts. - HANDOFF-codegen-coverage.md (new): self-contained handoff/resume prompt -- status, locations, the deterministic 4-process gate command set, open follow-ups, and the non-obvious gotchas (methodology-A gate, multiprocessing coverage jitter, driver-owns-gate, add-only, never-push). Distinct from HANDOFF.md (the prior #7 Validators effort). Docs only; no source, no test changes. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

…commit The previous implementation used currentBuild.previousBuild directly. On a "Restart from Stage", Jenkins creates a new build number, so previousBuild on build #7 points to build #6 (itself a restart), whose wfapi only contains stages from the restart point onward. After two restarts, the history of what passed in the original run is completely lost, causing all stages to re-run. Fix: walk the full previousBuild chain to collect all builds that ran on the same GIT_COMMIT as the current build, then query the wfapi of the EARLIEST build for that commit. The earliest build is the original (pre-restart) run which has the complete set of stage results, so the skip set correctly reflects everything that has ever passed for this commit regardless of how many restarts have occurred. Safety: if the commit changes, no prior builds match and the returned set is empty, so all stages run normally. Co-Authored-By: Claude Sonnet 4 <noreply@anthropic.com>

Adds _node_with_pos(node, capture) — combines _node_label (per-category- stream [N] index, plain MFMA omits) with bare '@ idx=M' position render. Single helper for the canonical Failure node reference shape, replacing three different prior styles: - #1, #2: manual `_node_label + " " + "@ idx=N"` concatenation - #4, #5: `category[name] format_position(...)` (rendered the FULL name inside brackets, e.g. `LRA0[LRA0[0]]`, plus a cross-category list suffix that duplicates [N]'s purpose) - #6: `category format_position(...)` (no brackets) All five formatters now route through _node_with_pos. Plain MFMA stays bracket-less per _node_label's MFMA discriminator; PackMFMAs (categories PackA*/PackB*) keep [N] because CMS reschedules them. Waits in #5 stay as bare '@ idx=N' — the surrounding 'SWaitCnt' word already names the kind; rendering the SYNC category as `SYNC[N] @ idx=M` would just duplicate that. Skipped: - #10 SCCConflict: brackets carry rocisa class name (e.g. [SCSelectB32]) not [N] index; semantic conflict, separate audit. - #7 WrongInterleaving / #8 TimingTooClose: use `name` field for Pack identity (MiddlePack_a/_b/_c, CVT0_a/_b); replacing with [N] would lose the a/b/c discriminator. - #13 ConstraintViolation: slated for deletion in bead `pcz`. 3 new pinning tests verifying [N] actually appears when capture is given (one per Failure: #4 LRA0[1], #5 LRA0[1] + plain MFMA bracket-less, #6 GRA[1] in trailing reference). Test count: 564 passed (+3).

…index Routes #7 WrongInterleavingFailure (pack/expected_next/actual_next) and #8 TimingTooCloseFailure (producer/consumer) through _node_with_pos for the canonical 'category[N] @ idx=M' rendering, dropping the per-formatter _idx() closures that read node.name + issued_at directly. Removes two redundant artifacts: 1. The duplicated vmfma_index in the prior message — GraphNode.name is `f"{category}@{vmfma}.{sub}"` so today's output renders as `PackA0@5.2 @ idx=5 has wrong interleaving...` (vmfma printed twice). Now: `PackA0[N] @ idx=5 has wrong interleaving...` (single source). 2. The dual-shape (`issued_at` vs `position`) detection. The structural- side ValidatorInstruction emitter was deleted in `ola.4` phase 2 (Pack rule); both Failures are now graph-side only. _node_with_pos's own getattr(node, 'position', None) or node.issued_at fallback handles any residual ValidatorInstruction caller without per-formatter closures. Tests rewritten: - test_wrong_interleaving_failure_format: drops the MiddlePack_a/b/c name-string assertions (vestigial — name was the OLD GraphNode label, no longer in the message). Pins the exact new rendering with empty capture. - test_timing_too_close_failure_format: same shape; pins exact new output. - 2 new tests added: test_wrong_interleaving_failure_format_with_capture_brackets and test_timing_too_close_failure_format_with_capture_brackets — verify [N] indexing kicks in when capture is real. Suite: 565 passed (was 563, +2 new bracket tests) / 2 skipped / 1 xfailed.

…rows in 8t9 audit DZL_SCOPE_REASSESSMENT.md §1.1 SCC: adds source-level confirmation that CommonInstruction (parent of every SCC-touching SOPx class) has only dst/dst1/srcs fields, with getDstParams()/getSrcParams() bodies shown verbatim from instruction.hpp:382. Rules out the q9j-style "binding gap" hypothesis at the source level, not just by Python dir() probing. ROCISA_WORKAROUNDS_AUDIT.md: adds a "Status update (2026-05-07)" preamble that flags which rows have been resolved or reframed by subsequent investigations: - Row #3 (_OPERAND_RULES, ~600 LoC estimate) → reassessed by q9j reassessment to ~12 lines C++ binding + ~260 LoC validator migration (the C++ API already exists). - Rows #7, #8, #19 (VCC machinery) → REMOVED by uraq; no rocisa-side replacement planned. VCC dataflow tracking is permanently dropped from the validator's scope. - MFMA acc special-casing → RECLASSIFIED to q9j Category A (acc and acc2 are already in the C++ getDstParams/getSrcParams partition). - Category C narrowed from SCC/VCC/m0/acc to SCC/m0 only. The audit's table and prose are preserved as the original artifact; the preamble lets a reader find the current canonical reframing for each affected row without rewriting the historical record.

… (rocm-libraries-wlrp) Adds Tensile/Components/DEFAULT_TO_CMS_CONVERTER_PHASE1.md scoping a function that takes a Tensile YAML, runs the default codegen, and produces a registered CMS schedule .py file equivalent to what the default emits. Headline finding — SNops are FULLY REPRESENTABLE in the CMS schedule format. The bead's 'if unrepresentable' contingency does NOT apply. 'SNOP' is a first-class category in the schema (CustomSchedule.py:222), ScheduleInfo.snopCode is a constructor field, the dispatcher walks optSchedule['SNOP'] identically to other categories, and existing schedule _get_schedule_96x256x64_16bit (CustomSchedule.py:1817+) already emits SNops successfully. Other findings: - CMS format mapped: ScheduleInfo(numCodePaths, numMfma, optSchedule, syncCode, nglshift, nllshift, nllZeroDscnt, mfmaReorder, snopCode). optSchedule = dict[str, list[list[int]]] — outer list per-codepath, inner list of mfma-slot positions. - Capture pipeline reuse: ctx.default.main_loop[0].instructions is list[TaggedInstruction] already pre-bucketed by category + slot.mfma_index. No new capture shape needed. - YAML loading: single entry point LibraryIO.read (LibraryIO.py:204) handles both default and alternate-format YAMLs. - Registered-schedule format: 3 imports + one @RegisterSchedule decorator. Import-time side effect via _SCHEDULE_REGISTRY.append. - Validator entry exists (compare_graphs at CMSValidator.py:2338, validate_edge_wait_coverage at :2751); no standalone YAML-driven CLI exists — this epic creates one. Critical formula correction (was wrong as written): nglshift / nllshift = (len(optSchedule['GRA'][0]) + len(optSchedule['GRB'][0])) // 2 NOT (len(GRA) + len(GRB)) // 2 — the latter would always return 1 for multi-codepath schedules. Audit against all existing schedules: 44 OK matches, 0 mismatches, 37 UNKNOWN-but-equivalent on inspection. No Phase 3 hazards. SNop emission fidelity (Phase 3 requirement): the dispatcher iterates optSchedule.items() in dict insertion order. The converter must populate keys in linear-stream order so the emitted instructions preserve the default-side order. Phase 1 established conclusions (no Phase 2 input needed): - E1: numCodePaths=1 policy. Default codegen yields a single linear stream; multi-codepath is human authoring optimization (out of scope per bead). - E2: nglshift formula (above) — validated against all existing schedules. - E3: TileConfig.isa is a free parameter; converter passes through the YAML's ISA (multi-arch confirmed). - E4: Default-side SNops are well-formed by construction (KernelWriter emits them via per-arch waitState calculations); a TimingTooCloseFailure from a default-side SNop would indicate a writer/profile bug, not a converter problem. 8 Phase 2 user decisions remain staged (down from 9 — #7 promoted to established). Most important: converter location, CLI shape, registration-collision policy. Bead stays open — Phase 2 needs user input before Phase 3 (TDD implementation) can begin.

…, #12, stacked on #1, #2] Add double-precision GEMM support to the CPU GEMM driver: - Add TypeTraits<double> mapping to DataType::Double - Add --type f64 branch calling runGemm<double, double> - Type all auxiliary buffers (bias, scaleAlphaVec, scaleAB) as AccumulateT so the slow path's GetValue(alphaType, ...) reads the correct byte width when alphaType=Double - Set alphaType/betaType to accumDtypeEnum so TypeId dispatch matches TypedGemm_D_D_D for the slow reference solver - Register bias/scale tensors with accumDtypeEnum - Add 10 f64 slow-path tests: 4 transpose combos, Beta, Bias, AllFeatures, TN_AllFeatures, ScaleAB Scalar/Vector Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

…, #12, stacked on #1, #2] Add double-precision GEMM support to the CPU GEMM driver: - Add TypeTraits<double> mapping to DataType::Double - Add --type f64 branch calling runGemm<double, double> - Type all auxiliary buffers (bias, scaleAlphaVec, scaleAB) as AccumulateT so the slow path's GetValue(alphaType, ...) reads the correct byte width when alphaType=Double - Set alphaType/betaType to accumDtypeEnum so TypeId dispatch matches TypedGemm_D_D_D for the slow reference solver - Register bias/scale tensors with accumDtypeEnum - Add 10 f64 slow-path tests: 4 transpose combos, Beta, Bias, AllFeatures, TN_AllFeatures, ScaleAB Scalar/Vector Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> (cherry picked from commit e12ae97)

…, #12, stacked on #1, #2] Add double-precision GEMM support to the CPU GEMM driver: - Add TypeTraits<double> mapping to DataType::Double - Add --type f64 branch calling runGemm<double, double> - Type all auxiliary buffers (bias, scaleAlphaVec, scaleAB) as AccumulateT so the slow path's GetValue(alphaType, ...) reads the correct byte width when alphaType=Double - Set alphaType/betaType to accumDtypeEnum so TypeId dispatch matches TypedGemm_D_D_D for the slow reference solver - Register bias/scale tensors with accumDtypeEnum - Add 10 f64 slow-path tests: 4 transpose combos, Beta, Bias, AllFeatures, TN_AllFeatures, ScaleAB Scalar/Vector Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

* test: add fp64 emulation unittest stub * build: use object libraries for linking * test: add small dgemm test with eager mode * build: don't require HIP to fp64 target * build: don't shadow library symbols * build: don't shadow library symbols 2

…ROCm#9 env var management)

…, #12, stacked on #1, #2] Add double-precision GEMM support to the CPU GEMM driver: - Add TypeTraits<double> mapping to DataType::Double - Add --type f64 branch calling runGemm<double, double> - Type all auxiliary buffers (bias, scaleAlphaVec, scaleAB) as AccumulateT so the slow path's GetValue(alphaType, ...) reads the correct byte width when alphaType=Double - Set alphaType/betaType to accumDtypeEnum so TypeId dispatch matches TypedGemm_D_D_D for the slow reference solver - Register bias/scale tensors with accumDtypeEnum - Add 10 f64 slow-path tests: 4 transpose combos, Beta, Bias, AllFeatures, TN_AllFeatures, ScaleAB Scalar/Vector Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> (cherry picked from commit e12ae97)

…, #12, stacked on #1, #2] Add double-precision GEMM support to the CPU GEMM driver: - Add TypeTraits<double> mapping to DataType::Double - Add --type f64 branch calling runGemm<double, double> - Type all auxiliary buffers (bias, scaleAlphaVec, scaleAB) as AccumulateT so the slow path's GetValue(alphaType, ...) reads the correct byte width when alphaType=Double - Set alphaType/betaType to accumDtypeEnum so TypeId dispatch matches TypedGemm_D_D_D for the slow reference solver - Register bias/scale tensors with accumDtypeEnum - Add 10 f64 slow-path tests: 4 transpose combos, Beta, Bias, AllFeatures, TN_AllFeatures, ScaleAB Scalar/Vector (cherry picked from commit e12ae97)

… uchar type (#9237) ## Motivation This PR addresses compile warnings with Ofast flag. Instead changed to use -fast-math ## Technical Details Modified compile flags to use standard O3 Removed the use of uchar data-type from RPP test-suite which is only defined in OpenCV ## Test Plan CTests should pass ## Test Result Start 1: rpp_sanity_test_brightness_host_f32 1/10 Test #1: rpp_sanity_test_brightness_host_f32 ......... Passed 3.64 sec Start 2: rpp_sanity_test_brightness_hip_f32 2/10 Test #2: rpp_sanity_test_brightness_hip_f32 .......... Passed 7.75 sec Start 3: rpp_qa_tests_tensor_image_host_all 3/10 Test #3: rpp_qa_tests_tensor_image_host_all .......... Passed 42.43 sec Start 4: rpp_qa_tests_tensor_voxel_host_all 4/10 Test #4: rpp_qa_tests_tensor_voxel_host_all .......... Passed 2.26 sec Start 5: rpp_qa_tests_tensor_misc_host_all 5/10 Test #5: rpp_qa_tests_tensor_misc_host_all ........... Passed 7.74 sec Start 6: rpp_qa_tests_tensor_misc_host_test_type_1 6/10 Test #6: rpp_qa_tests_tensor_misc_host_test_type_1 ... Passed 402.60 sec Start 7: rpp_qa_tests_tensor_image_hip_all 7/10 Test #7: rpp_qa_tests_tensor_image_hip_all ........... Passed 84.44 sec Start 8: rpp_qa_tests_tensor_voxel_hip_all 8/10 Test #8: rpp_qa_tests_tensor_voxel_hip_all ........... Passed 4.62 sec Start 9: rpp_qa_tests_tensor_misc_hip_all 9/10 Test #9: rpp_qa_tests_tensor_misc_hip_all ............ Passed 16.50 sec Start 10: rpp_qa_tests_tensor_misc_hip_test_type_1 10/10 Test #10: rpp_qa_tests_tensor_misc_hip_test_type_1 .... Passed 348.19 sec 100% tests passed, 0 tests failed out of 10 ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

…, #12, stacked on #1, #2] Add double-precision GEMM support to the CPU GEMM driver: - Add TypeTraits<double> mapping to DataType::Double - Add --type f64 branch calling runGemm<double, double> - Type all auxiliary buffers (bias, scaleAlphaVec, scaleAB) as AccumulateT so the slow path's GetValue(alphaType, ...) reads the correct byte width when alphaType=Double - Set alphaType/betaType to accumDtypeEnum so TypeId dispatch matches TypedGemm_D_D_D for the slow reference solver - Register bias/scale tensors with accumDtypeEnum - Add 10 f64 slow-path tests: 4 transpose combos, Beta, Bias, AllFeatures, TN_AllFeatures, ScaleAB Scalar/Vector (cherry picked from commit e12ae97)

* test: add fp64 emulation unittest stub * build: use object libraries for linking * test: add small dgemm test with eager mode * build: don't require HIP to fp64 target * build: don't shadow library symbols * build: don't shadow library symbols 2

jayhawk-commits and others added 6 commits April 30, 2025 18:46

Testing auto-label job post folder reorg

9e1efd0

Change grep logic in auto-labeler after local testing

e78b319

Testing autolabel with another change

94ea40a

Merge branch 'develop' into testAutoLabel

c170519

Merge remote-tracking branch 'origin/develop' into testAutoLabel

272f169

Merge branch 'testAutoLabel' of https://github.com/ROCm/rocm-libraries …

e0142fd

…into testAutoLabel

github-actions Bot added the project: rocthrust label Apr 30, 2025

jayhawk-commits closed this May 1, 2025

jayhawk-commits deleted the testAutoLabel branch May 9, 2025 23:10

jayhawk-commits self-assigned this May 17, 2025

jayhawk-commits added the migration Tasks or issues tied to migration to this monorepo label May 17, 2025

This was referenced Jun 30, 2025

Invalid instruction oops for gfx1151 on Fedora rawhide #458

Open

kernel oops for gfx1201 on Fedora Rawhide #459

Open

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

evbo mentioned this pull request Aug 14, 2025

[Feature] Tensile gfx1032 support #1202

Open

assistant-librarian Bot mentioned this pull request Sep 18, 2025

Kernrl oops on illegal instruction for gfx1151 on Fedora Rawhide #1675

Closed

ibrahimw1 mentioned this pull request Mar 12, 2026

[hipsolver] Refactor build system ibrahimw1/rocm-libraries#2

Open

1 task

WBobby mentioned this pull request Jun 18, 2026

[ALMIOPEN-2092] Fix selective rerun skipping catchError-failed parallel stages #8577

Merged

1 task

pghysels added a commit to pghysels/rocm-libraries that referenced this pull request Jun 19, 2026

merge: integrate braden/fp64_emulation_develop (ROCm#7 unit test stub, …

13fdbc1

…ROCm#9 env var management)

BrianHarrisonAMD mentioned this pull request Jun 26, 2026

[Issue]: miopen_plugin ctest runs duplicate integration-test stages (quick≡standard, comprehensive≡full) — 4× redundant runs that amplify flakiness and waste CI time #8852

Open

rrawther mentioned this pull request Jul 9, 2026

fix(rpp): Remove Ofast compiling option from CMakeList.txt and change uchar type #9237

Merged

1 task

talumbau mentioned this pull request Jul 16, 2026

feat(hipblaslt): GEKO — GEMM Kernel Optimizer #8832

Merged

luraess mentioned this pull request Aug 12, 2026

Segfaults in ConvHipImplicitGemmGroup{Fwd,Wrw,Bwd}Xdlops on gfx942 (MI300) for small/degenerate convolutions — MIOpen 7.2.3 #9088

Open

Navigation Menu

Uh oh!

Testing auto-label job post folder reorg - #7

Testing auto-label job post folder reorg#7
jayhawk-commits wants to merge 6 commits into
developfrom
testAutoLabel

jayhawk-commits commented Apr 30, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Uh oh!

Conversation

jayhawk-commits commented Apr 30, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant