[hipDNN] Resolve conflicts with latest ROCm/rocm-libraries develop. - #1
Closed
CMiservaAMD wants to merge 236 commits into
Closed
CMiservaAMD wants to merge 236 commits into
CMiservaAMD wants to merge 236 commits into
Conversation
## Motivation Adds `build` API to easily call all the following methods: ```cpp HIPDNN_CHECK_ERROR(validate()); HIPDNN_CHECK_ERROR(build_operation_graph(handle)); HIPDNN_CHECK_ERROR(create_execution_plans(modes)); HIPDNN_CHECK_ERROR(check_support()); HIPDNN_CHECK_ERROR(build_plans()); ``` ## Technical Details Adds the `build` method to `Graph`, replaces legacy call sequence with `build` in samples and integration tests, and adds tests. ## Test Plan Run tests & samples. ## Test Result Tests pass, and a few samples have unrelated failures. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
## Motivation Upgrade rocFFT and hipFFT to use the C++20 standard. This modernization enables access to newer language features, improved compile-time diagnostics, and aligns with the broader ROCm ecosystem's move toward C++20. ## Technical Details - Updated runtime compilation flags from -std=c++14 to -std=c++20 in rtc_compile.cpp - Updated external Boost build flags from -std=c++11 to -std=c++20 * Refactored Statement in generator.h from a type alias to a class wrapper around std::variant to satisfy C++20's stricter requirements for complete types * Fixed deprecated implicit this capture in lambdas ([=] → [=, this]) in rocfft_ostream.cpp, stockham_gen_base.h, plan.cpp, and hipfftw_helper.h * Fixed std::accumulate lambda signatures in hipfftw_test.cpp and gtest_main.cpp to use pass-by-value instead of lvalue references, as required by C++20 * Wrapped unreachable code in hipfft_accuracy_test.cpp with #ifdef _CUFFT_BACKEND to eliminate warnings ## Test Plan Check full compilation on different systems, verify correctness and no regression. ## Test Result All rocff-test and hipfft-test passed on Linux. Complete CI in progress. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Steve Leung <Steve.Leung@amd.com>
## Motivation Remove use of Boost scope exit. There's no really good reason to pull in Boost just to save a couple of lines of uncomplicated code. ## Test Plan Run existing tests. ## Test Result Tests pass. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
## Motivation
Our previous configuration system had become limiting in several ways.
Most importantly, it was not able to differentiate between individual
GPUs when selecting config parameters. This made proper tuning difficult
and prevented future work involving SPIR-V–specific tuning. In addition,
the old approach relied heavily on complex template metaprogramming,
which had become difficult to maintain. With the move to C++17, we now
have cleaner and more expressive language features available, making
this a good opportunity to redesign the system.
## Technical Details
All changes are internal. **There are no API changes for users.**
The majority of the diff in this PR consists of the new configuration
definitions themselves, so while the PR appears large, the actual code
changes are relatively small.
### New Configuration Structure
Each algorithm now defines a *_config_picker templated on the target and
value type. Below is a simplified example:
```cpp
template<class Target, class value_type>
constexpr <algo_name>_config_picker()
-> std::enable_if_t<
std::is_same_v<Target,
comp_target<gen::gcn5, target_arch::gfx906, gpu::mi50, rep::amdgcn>>,
<algo_name>_config_params>
{
// Tuned configuration #1
if constexpr (/* condition for this combination */)
{
return <algo_name>_config_params{ ... };
}
// Tuned configuration #2
if constexpr (/* condition for this combination */)
{
return <algo_name>_config_params{ ... };
}
// Default for this target
return <algo_name>_config_params_base<value_type>();
}
```
Each tuned target provides a similar overload. For untuned or unknown
targets, we provide a general fallback:
```cpp
template<class Target, class value_type>
constexpr auto <algo_name>_config_picker()
-> std::enable_if_t<
std::is_same_v<Target,
comp_target<gen::unknown, target_arch::unknown, gpu::generic, rep::amdgcn>>,
<algo_name>_config_params>
{
// Fallback: use a commonly tuned target (often MI100)
return <algo_name>_config_picker<
comp_target<gen::cdna1, target_arch::gfx908, gpu::mi100, rep::amdgcn>,
key_type, value_type>();
}
```
All available tuned targets are listed in:
```cpp
using <algo_name>_targets = comp_targets<
comp_target<gen::gcn5, target_arch::gfx906, gpu::mi50, rep::amdgcn>,
...,
comp_target<gen::unknown, target_arch::unknown, gpu::generic, rep::amdgcn>>;
```
### How Config Selection Works Now
In the new system, kernels are compiled for all tuned targets. At
runtime, if the current GPU does not have dedicated tuning, the library
uses the most_common_config policy to choose the best matching compiled
kernel.
The selection policy (tested in test_config_dispatch.cpp) attempts to
match, in decreasing priority:
1. Exact GPU model
2. Architecture
3. Generation
If no match is found, it falls back to the unknown target. If multiple
candidates match, the last one listed in the comp_targets type list is
chosen, which gives us a controlled and predictable fallback order.
We also pass the selected target into kernel compilation, enabling
compile-time specialization based on GPU, architecture, and generation.
### Target struct
The target struct currently stores only:
- GPU generation
- Architecture
- GPU Name
- Representation (rep), which distinguishes SPIR-V from native AMDGCN
The rep field is not yet functional (requires compiler support), and the
dispatch policy does not consider it at the moment. Also this target
structs makes it relatively easy to store more data.
### Scripts
The python script changes in this PR are there for scripts that used the
configs as input/output.
### Summary of Improvements:
- Better differentiation and selection across GPUs
- Cleaner C++17-based implementation
- Easier extension for future SPIR-V tuning
- Improved maintainability of config definitions
- Added more flexibility for future features.
## Test Plan
Some tests were added in test_config_dispatch.cpp, these and all the
other tests should pass. Also everything needs to be benchmarked to see
if the correct configs are chosen.
## Test Result
All tests pass, benchmarks are still WIP.
## Submission Checklist
- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
---------
Co-authored-by: Saiyang Zhang <saiyang@streamhpc.com>
* Support for gfx115X * Changes for gfx115X * Add gfx1153 * Update changelog --------- Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
## Motivation Adds herk_ex which was added to rocBLAS. ## Technical Details Follow standard pattern for rocBLAS additional API exposure Fixes syrk_ex reference to support bfloat16 output ## Test Plan Update yaml test sets for standard harness. Extensive functionality testing in rocBLAS is not replicated at this library level. Added smoke tests for syrk_ex and herk_ex.
…no longer needed (#3603) ## Motivation Now that we have completely customizable tests for each backend (ROCm/CUDA), we do not need the platform checks inside the testing_xxx.hpp files. This PR removed them to simplify and shorten the test code.
## Motivation Enable gfx1150, gfx1152, and gfx1153 targets. ## Technical Details Add these targets to DEFAULT_AMDGPU_TARGETS in hipCUB CMakeLists.txt. ## Test Plan Build existing ctests for, and run them on, the new targets. ## Test Result - [x] gfx1150 passed - [x] gfx1152 passed - [x] gfx1153 passed ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
* [CK_TILE] unify double and single lds implementation (#108) Unify LDS buffer management API for single and double buffering modes This change consolidates the Local Data Store (LDS) buffer management by: Merging single and double LDS buffer APIs into a unified interface Implementing ping-pong address calculation in pipeline when double LDS is enabled Computing pong buffer addresses dynamically using base address offsets --------- Co-authored-by: joye <joye@amd.com> Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * update wp_pipeline * fix a c++17 issue * update for ci errors * fix ci issues * include a header to fix ci errors * fix some rebase issues * update with rebase --------- Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
…nation for channel-only tensors (#3608) ## Motivation Samples are failing verification for the following reasons: 1. The batch norm applicability checks discriminate between NCHW and NHWC strides for tensors with channel only dims. However, the strides are functionally identical since they reference the same contiguous block of C elements in memory. The samples generate NCHW strides universally via omitting the layout specifier, and they failed when they hit these applicability checks. 3. The integration tests for BN fwd training uses `getRmsToleranceTraining` with `allClose`, whereas the samples use `allClose` with the regular `getToleranceTraining`. We should be using `getToleranceTraining` with a suitable tolerance with `allClose`, and reserve `getRmsToleranceTraining` for use with the RMS validator. ## Technical Details - Update the applicability checks to accept channel-only tensors with any layouts - Swap the use of `getRmsToleranceTraining` with `getToleranceTraining` in the integration tests. - Update BN training tolerances to be the minimal values that allow our tests and samples to pass. ## Test Plan Run tests & samples. ## Test Result Tests & samples pass. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
## Motivation Fail the test when the kernel launch fails. This is relevant to detect obvious problems in test setup, for example using unit test binaries that were build for a different GPU arch. Adds `CHECK_HIP_ERROR(hipGetLastError());` after every call to `hipExtLaunchKernelGGL` in the rocWMMA unit tests. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
## Motivation Update some kernels in Origami+SK GFX942 BBS libs. ## Technical Details These kernels give better performance. ## Test Plan CI and Locally. ## Test Result Improved Performance. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
The dyna mpi worker compilation had been failing due to a missing cmake dependency; this fixes that issue.
## Motivation hipDNN is intended to remain host-only, and be a generic target for TheRock. Since this plugin test had a kernel inside it, hipDNN had a gpu target requirement. Now that we have plugins for hipDNN this test path is exercised elsewhere, and it's no longer needed. Fixes ROCm/TheRock#2758 ## Technical Details Remove test plugin that required kernel code to execute. ## Test Plan Run tests, tidy, and code coverage to ensure everything is working correctly. ## Test Result Tests run correctly.
* fix some issues from internal branch * update cshuffle_epilogue * update cshuffle_epilogue * update cshuffle * update warp_gemm
* ck-builder: explicitly delete forward declarations Before, these functions were seen as a forward declaration for an existing function. If no actual implementation overload could be found, these would be selected and a linker error or warning would be generated. By marking these functions as explicitly deleted, they incorrect invocations are generated as compile error instead. * ck-builder: ckt::run plumbing for reference conv This implements the ckt::run plumbing for the reference convolution implementation and sets up the first complete end-to-end test. * ck-builder: make validation system check for all-zeros When both the actual and reference output are both all zero bits, there is probably something wrong in the test framework. * ck-builder: proper implementation+tests for TensorDescriptor::is_packed * ck-builder: fix typos
## Motivation Several unit tests were failing when run on therock on MI200. These tests failed because they relied on CK, and CK is currently not built on some older targets by default (in therock). ## Technical Details The fix is to skip the tests that require CK if CK is not built. ## Test Plan Build without CK and then run the following tests on MI200: - bin/miopen_gtest --gtest_filter="Smoke/GPU_ConvBiasResAddActivation_fwd_*" - bin/miopen_gtest --gtest_filter="Smoke/GPU_KernelTuningNetTestConvHipIgemmGroup*" - bin/miopen_gtest --gtest_filter="GPU_FusionPlan_FP16*" The above tests cover all the failures that were seen. ## Test Result All the tests pass ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
## Motivation Increase architecture support in hipTensor by supporting: - gfx1100 - gfx1101 - gfx1102 - gfx1103 - gfx1150 - gfx1151 - gfx1152 - gfx1153 ## Technical Details - The architectures support up to 32 bits data type, and 16 bits matrix core compute type (with WMMA). - gfx115X support depends on CK's ROCm/composable_kernel#3496 ## Test Plan - Manually run all the tests in the newly supported architectures. ## Test Result - gfx1100: Pass - gfx1101: Pass - gfx1102: Pass - gfx1103: Not executed - gfx1150: Pass - gfx1151: Pass (Linux kernel 6.14+ is required - https://rocm.docs.amd.com/projects/radeon-ryzen/en/latest/docs/install/installryz/native_linux/install-ryzen.html#prepare-the-system) - gfx1152: Pass - gfx1153: Pass ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
- These are triggering a compiler error with newer compiler
…tead (#3627) ## Motivation Alternative approach to #3591. ## Technical Details Keep the user-provided mpi launch command as a `std::string` option to rocfft-test and use an internal `CLI::App` object to parse it thereafter to extract the command's `argv` from it. ## Test Plan Multi-processes and/or multi-device tests cover the changes. ## Test Result Tests pass. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
…l for' (#3642) ## Motivation and details The current usage is incorrect as it forces the interruption of the parallel loop _before_ the value to be returned was effectively updated by the thread finding a conflict when/if the `cancel` constructs are actually enabled (which is [not necessarily the default](https://www.openmp.org/spec-html/5.0/openmpse59.html)): without this suggested fix, tests may fail when/if the environment variable `OMP_CANCELLATION` is set to `true`. ## Test Plan The [`valid_length_stride`](https://github.com/ROCm/rocm-libraries/blob/9b2657993554a1345e46c70a2898b88bc8b692cd/projects/rocfft/clients/tests/validate_length_stride.cpp#L124) tests were executed manually with the environment variable `OMP_CANCELLATION` defined and set to `true`. ## Test Result All tests pass with the suggested changes (some fail otherwise). ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
hipblaslt was not fully excluded from pre-commit. The proper path here would be to raise a PR that fully formats components, and then simultaneously turns on pre-commit linting/formatting for that section of code. This prevents "piecemeal" linting/formatting. We should turn on these pre-commits soon (and then re-enable the noted exclusions).
…#3648) ## Motivation <!-- Explain the purpose of this PR and the goals it aims to achieve. --> Add note explaining rocBLAS dependency. ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> Modify install guides ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> Build locally ## Test Result <!-- Briefly summarize test outcomes. --> NA ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
github-actions
Bot
added
project: hipcub
project: hipfft
project: hipsolver
project: hipsparse
project: hipsparselt
project: miopen
project: rocblas
project: rocfft
project: rocprim
project: rocsolver
project: rocsparse
project: rocthrust
shared: mxdatagenerator
shared: rocroller
github actions
labels
Saiyang-Zhang
pushed a commit
that referenced
this pull request
Feb 2, 2026
ROCm#3710) ## Motivation Optimizing the tensor filling functions started a discussion about optimizing tensor iteration in general: ROCm#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 ROCm#2: hipdnn_backend_tests ............. Passed 1.29 sec Start 3: hipdnn_frontend_tests 3/7 Test ROCm#3: hipdnn_frontend_tests ............ Passed 0.03 sec Start 4: hipdnn_test_sdk_tests 4/7 Test ROCm#4: hipdnn_test_sdk_tests ............ Passed 4.32 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.33 sec Start 7: public_hipdnn_frontend_tests 7/7 Test ROCm#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>
matyas-streamhpc
pushed a commit
that referenced
this pull request
Mar 5, 2026
## Motivation <!-- Explain the purpose of this PR and the goals it aims to achieve. --> This PR addresses hipDNN issue ROCm#4951, which requests adding missing frontend integration test coverage for the Matmul operation. ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> 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 <!-- Explain any relevant testing done to verify this PR. --> All tests were built and executed inside the official TheRock docker environment. ## Test Result <!-- Briefly summarize test outcomes. --> ``` [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 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>
matyas-streamhpc
pushed a commit
that referenced
this pull request
Mar 9, 2026
…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 #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 #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 #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.
ex-rzr
pushed a commit
that referenced
this pull request
Mar 27, 2026
…Cm#5880) ## Motivation Fix a bug in the smart-build --ctest-only filter that was incorrectly excluding tests with numbers less than 100. ## Technical Details The issue was caused by CTest formatting test numbers with variable spacing based on the number of digits: - "Test `#1`: name (3 spaces for tests 1-9)" - "Test `ROCm#79`: name (2 spaces for tests 10-99)" - "Test `ROCm#100`: name (1 space for tests 100+)" The previous code used `line.strip().startswith("Test #")` which only matched tests with a single space (i.e., test numbers >= 100). This caused tests like ck_tile_unit_sequence (Test ROCm#79) to be excluded from smart-build test selection, resulting in CTest failures when the binary wasn't built. Solution: Replace string matching with a regex pattern that handles all spacing variations: r'^\s*Test\s+#\d+:\s*(.+)$' ## Test Plan Tested with test numbers from 1 to 12345. ## Test Result - Before: 48 tests selected (only tests ROCm#100+) - After: 146 tests selected (all CTest-registered tests) ## Submission Checklist - [x ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com> Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
cenxuantian
pushed a commit
that referenced
this pull request
May 4, 2026
) # Add gfx950 MXFP4 Subtile-based kernel implementation ## Summary This PR is a follow-up to ROCm#6499 ([hipblaslt] Add support for gfx950 mxfp4) and adds the **Subtile-based kernel implementation (`UseSubtileImpl=1`)** for hipBLASLt on **gfx950**. It introduces a new tile-decomposed code generation path optimized for **MXFP4** and **BF16** GEMMs, plus the solution-selection plumbing, validation, Origami logic yamls, and unit tests needed to make it production-usable. ## Motivation PR ROCm#6499 brought MX data type support online for gfx950, but the existing TensileLite codegen path leaves significant performance on the table for MXFP4-heavy workloads. The Subtile path restructures global-read / local-read / MFMA / store scheduling at a finer granularity, which **greatly improves MXFP4 GEMM performance when using `HIPBLASLT_MATMUL_MATRIX_SCALE_BLK32_UE8M0_32_8_EXT`** (added to the hipBLASLt CHANGELOG). ## What's included ### 1. New Subtile-based kernel components (Tensile) New modules under `projects/hipblaslt/tensilelite/Tensile/Components/`: * `SubtileBasedKernel.py` (~1850 LOC) — entry point and orchestration of the subtile codegen path; replaces large portions of the standard prefetch / unroll / store flow when `UseSubtileImpl=1`. * `SubtileBasedLogicalScheduler.py` (~2415 LOC) — logical scheduler that builds the subtile-grained instruction graph (GR loads, LR offsets, MFMA tiles, scale loads, stores) from kernel parameters. * `SubtileBasedInstructionScheduler.py` (~433 LOC) — converts the logical schedule to an emit order respecting wave / register / hazard constraints. * `SubtileBasedInstructionEmitter.py` (~216 LOC) — instruction emission helpers shared by the subtile components. ### 2. Kernel writer / common changes * **`KernelWriter.py`**, **`KernelWriterAssembly.py`**: integration points for the subtile path — prefetch, GR offset calculation, LR offset calculation, post-loop, MFMA macro accounting, optimized `storeD`, LDS buffer swap, MX FP4 scale emit, `SrdMXSA/B+2` handling, sgpr allocation / overflow guards, computeLoadSrd fix. * **`SolutionStructs/Solution.py`**, **`SolutionStructs/Problem.py`**: introduces the `UseSubtileImpl` parameter, MX-related reject conditions for non-Subtile paths on gfx950, and additional valid GEMM type combinations for MX inputs. * **`Common/ValidParameters.py`**, **`Common/RequiredParameters.py`**, **`Common/GlobalParameters.py`**: `UseSubtileImpl` registration and defaults. * **`Components/StreamK.py`**: subtile-aware StreamK fixup (incl. import union with the `BufferLoadB32` cache-coherence change from ROCm#6837). * **`Components/GlobalWriteBatch.py`**: optimized global write batching for the subtile path (~670 LOC of changes). * **`Components/ComputeStoreVgprs.py`**, **`Components/LSU.py`**, **`Components/WorkGroupMappingAlgos.py`**, **`AsmStoreState.py`**, **`KernelWriterModules.py`**: minor adjustments needed by the subtile pipeline. ### 3. rocisa / host / client * **`rocisa/rocisa/include/container.hpp`**: helpers needed by the new emitter. * **`tensile_host.cpp`**, **`include/Tensile/TensorDescriptor.hpp`**: small fixups for the subtile path and gfx950 build. * **`client/include/DataInitialization.hpp`**, **`client/src/DataInitialization.cpp`**, **`client/src/Reference.cpp`**, **`client/src/ReferenceValidator.cpp`**, **`client/include/TypedId.hpp`**: MX scale init and reference paths used by the new tests. * **`clients/common/include/testing_matmul.hpp`**, **`clients/common/include/norm.hpp`**, **`clients/common/include/hipblaslt_datatype2string.hpp`**, **`clients/common/src/mxDataGen.cpp`**: wiring for batched (>1) testing and MX init. ### 4. Origami / solution selection (gfx950 MXFP4) New auto-tuned logic yamls under `projects/hipblaslt/library/.../Tensile/Logic/asm_full/gfx950/gfx950/Origami/` covering the FP4 SS / HS / BS variants in three layouts: * `Origami/` (default) * `Origami/Origami_nta4/` (no-transpose-A FP4) * `Origami/Origami_ntb4/` (no-transpose-B FP4) (9 new `gfx950_Cijk_Alik_Bljk_F4{SS,HS,BS}_MXA32_MXB32_*_UserArgs.yaml` files in total.) ### 5. New tests **End-to-end gfx950 GEMM yamls** in `Tensile/Tests/common/gemm/gfx950/`: * `subtile_bf16.yaml`, `subtile_mxfp4.yaml` * `mx32f4_tn.yaml`, `mx32f8_tn.yaml` * `mxfp4_mxfp4_{fp32,bf16}_tn_act{,_groupgemm}.yaml` * `mxfp4_fp8_{fp32,bf16}_tn_act{,_groupgemm}.yaml` * `fp8_mxfp4_{fp32,bf16}_tn_act{,_groupgemm}.yaml` **StreamK + MX:** `Tensile/Tests/common/streamk/sk_mx32f4_quick.yaml`, `sk_mx32f8_quick.yaml`. **New unit tests** (`Tensile/Tests/unit/`): * `test_SubtileBasedLogicalScheduler.py` (~1735 LOC) * `test_SubtileBasedSchedulerRef.py` (~596 LOC) * `test_gr_lr_roundtrip.py` (~571 LOC) * `test_storeD_roundtrip.py` (~2420 LOC) * `test_graTileAssignment.py` (~354 LOC) * `test_lraTileAssignment.py` (~360 LOC) * `conftest.py`, `gpu_test_helpers.py` shared fixtures (~601 LOC) **New gtest:** `tensilelite/tests/MXScalePadding_test.cpp`. ### 6. Misc / hardening * Reject conditions: gfx950 MX + non-Subtile, DepthU constraints, GroupGEMM not yet supported with StreamK + MX, AssertSummationElementMultiple=256 for subtile MXFP4, missing-mxblock check for non-MX types. * Skip rocRoller for FP4-A/FP4-B with pre-swizzled scale layout (ROCm#42). * `forceDenorm=False` in `generateMXInput` (ROCm#11). * Several rebase fixes, copyright/year header updates, and review-comment fixes to `KernelWriter` / `KernelWriterAssembly`. ### 7. CHANGELOG Greatly improved MXFP4 GEMM performance when using HIPBLASLT_MATMUL_MATRIX_SCALE_BLK32_UE8M0_32_8_EXT ## How to use Set `UseSubtileImpl: 1` on a gfx950 MX-FP4 solution (see the new `subtile_mxfp4.yaml` / `mx32f4_tn.yaml` for canonical configs). The path is opt-in — non-MX and non-gfx950 kernels are unaffected. ## Backwards compatibility / risk * All new behavior is gated on `UseSubtileImpl=1` and gfx950. Existing solutions on other architectures or non-MX paths are unchanged. * `GroupGEMM + StreamK + MX` is intentionally rejected for now (TODO). * New Origami yamls only add solutions; nothing existing is modified. ## Test plan * New gtests + unit tests run automatically in CI (Tensilelite Python unit suite, `MXDataGen_test`, `MXScalePadding_test`). * New end-to-end gfx950 GEMM and StreamK yamls are added to the common test buckets. * Manual: run the gfx950 MXFP4 subtile suites (`pytest -k gfx950` after building Tensile, plus `tensilelite-client --yaml subtile_mxfp4.yaml` for sanity). ## Notes for reviewers * This branch was rebased onto current `develop` (post-ROCm#6499) by skipping the `users/nakajee/gfx950_mx_rebase_merge` history (which ROCm#6499 squash-merged) and replaying only the subtile-specific work as a single squashed commit. The actual code changes in this PR are limited to the files listed above (24 added, 56 modified; ~+170k / −2.6k including generated logic yamls). * The largest reviewable diffs are: * `Tensile/Components/SubtileBased{Kernel,LogicalScheduler,InstructionScheduler,InstructionEmitter}.py` (new files) * `Tensile/KernelWriter.py`, `Tensile/KernelWriterAssembly.py` * `Tensile/SolutionStructs/{Problem,Solution}.py` * `Tensile/Components/{GlobalWriteBatch,StreamK}.py` * `clients/common/include/testing_matmul.hpp` * `client/src/DataInitialization.cpp` * Description of all commits that were squashed for this feature branch: Subtile implementation for gfx950 MX FP4 --- 272f88d: Add sample subtile impl --- Author: brianshi <brianshi@amd.com> --- 60ecede: GR Offset calculation (#1) --- Author: sebvince <115461989+sebvince@users.noreply.github.com> --- be69c1d: Enable post-loop code generation, and add some subroutines --- Author: b-shi <brianshi@amd.com> --- 646d102: LR offset calculation (ROCm#2) --- Author: sebvince <115461989+sebvince@users.noreply.github.com> --- 71f4bca: Add GR load emit logic, and misc fixes (ROCm#3) --- Author: b-shi <brianshi@amd.com> --- 1fd0db9: Emit LR + init ACCVGPR (ROCm#4) --- Author: sebvince <115461989+sebvince@users.noreply.github.com> --- 9d406b9: Add loop and ptr update code --- Author: b-shi <brianshi@amd.com> --- b6127bc: Update GR/LR offset calculation to fully support 2x2, 1x4, 4x1 waveConfigs (ROCm#7) --- Author: sebvince <115461989+sebvince@users.noreply.github.com> --- 89ec87c: Account for valuC macro value in SK WS store code --- Author: b-shi <brianshi@amd.com> --- 6edf53d: Rebase fix --- Author: b-shi <brianshi@amd.com> --- 34e79fc: Enable fp4 (ROCm#8) --- Author: sebvince <115461989+sebvince@users.noreply.github.com> --- d5a5c57: [Tensilelite] Add MX FP4 scale offset computation for subtile-based kernel (ROCm#6) --- Author: Archana Ramalingam <98564406+archana-ramalingam@users.noreply.github.com> --- 7a8a85a: Add lds buffer swap logic --- Author: b-shi <brianshi@amd.com> --- d24a8fe: Add optimized storeD code (ROCm#9) --- Author: b-shi <brianshi@amd.com> --- a45c20c: Fix MX scale tensor initialization: set forceDenorm=false in generateMXInput (ROCm#11) --- Author: T.J. Alumbaugh <T.J.Alumbaugh@amd.com> --- f945268: [Tensilelite] Enable the MX FP4 scale emit code in the subtile-based kernel (ROCm#10) --- Author: Archana Ramalingam <98564406+archana-ramalingam@users.noreply.github.com> --- cf37df4: Use fixed value for SrdMXSA/B+2 (ROCm#14) --- Author: Koji Nakajima <75698246+nakajee@users.noreply.github.com> --- f0c8dbc: Merge subtile_mx_f4_schedule to subtile_mx branch (ROCm#16) --- Author: b-shi <brianshi@amd.com> --- 543796f: Enable DU > 256, and reduce sgpr allocation (ROCm#18) --- Author: b-shi <brianshi@amd.com> --- c65bdb0: Add missing mxblock check for non-mx data types --- Author: b-shi <brianshi@amd.com> --- d64d226: Introduce UseSubtileImpl parameter (ROCm#20) --- Author: b-shi <brianshi@amd.com> Squash commits 20-35 from subtile_mx branch --- e4780da: Enable FixSrd2 for A/B (ROCm#23) --- Author: b-shi <brianshi@amd.com> * Enable FixSrd2 for A/B * Address comments from PR --------- --- e4c64a7: Add nt libs --- Author: b-shi <brianshi@amd.com> --- cd13ec1: [Tensilelite] Pad MX scale tensor dimensions for unaligned problem sizes (ROCm#21) --- Author: Archana Ramalingam <98564406+archana-ramalingam@users.noreply.github.com> * Add scale padding * Add tests * Remove redundant pre-swizzle path * Remove code from conflict * Fix reverted mxdatagen path for tensile tests * Add diverse test cases for scale padding in MXScalePadding_test and subtile.yaml - Expanded test cases to include non-multiple-of-32, even non-multiple-of-16, and odd dimensions. --- d87938f: Split subtile.yaml into subtile_bf16.yaml and subtile_mxfp4.yaml (ROCm#22) --- Author: James Newling <james.newling@gmail.com> Replace the 'monolithic' subtile.yaml with two focused test files. All original test coverage is preserved. Two new FP4 groups added. BF16 coverage (subtile_bf16.yaml, tests are essentially unchanged): # | Description | Dest | MIs | PGR | DU | SK | Sizes --+--------------------+------+-----+-----+---------+-----+------ 0 | BF16 TN main | b | 19 | 0 | 64 | 0,3 | 11 1 | BF16 TN large DU | b | 4 | 0 | 128,192 | 0,3 | 7 2 | BSS (f32 output) | s | 6 | 0 | 64 | 0,3 | 9 3 | BF16 bias | b | 2 | 0 | 64 | 0 | 1 FP4 coverage (subtile_mxfp4.yaml): # | Description | Dest | MIs | PGR | DU | SK | Sizes | Status --+--------------------+------+-----+-----+-----+-----+-------+-------- 0 | FP4 TN main | b | 15 | 0 | 256 | 0,3 | 23 | from original 1 | FP4 TN large DU | b | 4 | 0 | 512 | 0,3 | 13 | from original 2 | F4SS (f32 output) | s | 5 | 0 | 256 | 0,3 | 13 | from original 3 | FP4 bias | b | 2 | 0 | 256 | 0 | 1 | from original 4 | FP4 PGR=2 | b | 13 | 2 | 256 | 0 | 5 | new 5 | FP4 expanded MIWT | b | 24 | 0 | 256 | 0 | 5 | new 6 | PGR=2 WG 4x1/1x4 | | 6 | 2 | 256 | 0 | 1 | known failures (commented) Run times on gfx950 (8x MI350X): File | NEV=-1 | NEV=0 -------------------+--------+------ subtile_bf16.yaml | 23s | 23s subtile_mxfp4.yaml | 37s | 40s Where NEV is number of elements to validate. I (James) have checked these numbers, and weirdly it is true that NEV=0 is a bit faster than NEV=-1 for mxfp4. --- af04f0d: Dependency based instruction scheduling (ROCm#19) --- Author: sebvince <115461989+sebvince@users.noreply.github.com> * Revert to single partition * Start using dependencies * as is * start using separate EmittedModules * remove reduntant wait * Add _extractPathsFromBeforeDeps * Continue simplification * Simplifying * Add more rules * cleanup * Add fp4 test * fix test * Add tests * Remove after field on emittedmodule * Refactoring instructionSchedule * Add comments * cleanup modules vs ops * Refactoring print functions * Test cleanup * Add more tests * Replace subgroup by partition * Remove unused unroll param * Add high level notes * Simplify NLL and NGLL GR removal * Add some comments * Force instruction insertion if no slots available * Fix test after rebase * Move scale before A/B and track inflight count * Fine-grain vmcnt calculation * Separate counts for scaleA and B * Avoid using m0 update and buffer_lod on same MFMA slot to avoid scalar instruction serialization * Fix test * Add vmcnt test * Fix duplicated loads for 1x4 and 4x1 * Fix placement in reverse order * Fix regression on PGR0 * add fallback to numMFMA=1 --- 3ec902b: Add some 1x4 and 4x1 origami solutions --- Author: b-shi <brianshi@amd.com> --- c5000d3: Fix typo --- Author: b-shi <brianshi@amd.com> --- 226ed84: [hipblaslt] Refactor Srd2 calculation for useFixedSrd2 (ROCm#30) --- Author: Koji Nakajima <75698246+nakajee@users.noreply.github.com> --- abf19d4: [Tensilelite] UseSubtileImpl: subtile-aligned edge check for store path (ROCm#29) --- Author: b-shi <brianshi@amd.com> * [Tensilelite] UseSubtileImpl: subtile-aligned edge check, OOB guard, and refactoring - Replace Size%MT edge check with subtile-aligned check: NonEdge paired store when trailing rows/cols are a multiple of the subtile block size (waveGroupM rows for M, 16 cols for N). Non-last workgroups always take NonEdge. - Add per-wave OOB guard (subtileM32ValidBlocksSgpr / subtileN16ValidBlocksSgpr) to skip stores outside valid M/N tile bounds in the NonEdge path. - Refactor duplicated OOB guard into _emitSubtileOobGuard helper; refactor M/N guard SGPR computation into _emitSubtileMGuard / _emitSubtileNGuard. - Fix orphan scalar store blockIdxM (was tt0, now (tt0*MatrixInstM)//mBlockSize). - Add quick-exit and edge/non-edge header comments to generated ASM. * Add some bias tests, combine M/N guard to single routine * Add OOB check for C loads, update storeD unit tests to check OOB, simplify quick exit checks * Address more PR comments: add M group skip, and skip to store end. simplified loadC OOB mask --------- --- 637881a: Fix unit tests & remove legacy code for subtile interleaving (ROCm#33) --- Author: sebvince <115461989+sebvince@users.noreply.github.com> * Fix gr_lr_roundtrip test * Use non-interleaved version as ref code * Fix scheduler test * Removed legacy interleaved mode for LR/GR offset calculation --- e9cb889: Fix MX FP4 scale buffer allocation and initialization for batched GEMM (ROCm#25) --- Author: Archana Ramalingam <98564406+archana-ramalingam@users.noreply.github.com> * Fix bacth count issue * Add batch count tests * Fix bacth count issue * Address PR review: clarify FP4-specific byte stride and add non-aligned batched tests - Updated comments on dataBatchBytes computation to clarify FP4 packing assumption (2 elements/byte) and flag that non-FP4 block-scaling types would require updating this conversion. - Added batched test cases with non-multiple-of-32 M/N dimensions: FP4 DU=256: [48,48,2] and [33,65,2] FP4 DU=512: [63,63,2] BF16: [50,100,2] --------- --- a43247b: Update some test yamls (ROCm#31) --- Author: b-shi <brianshi@amd.com> --- e2f69c8: Add f4bs origami library with activation function support. Refactor sgpr allocation to reduce sgpr usage in post loop. Store code-path reorganization (ROCm#32) --- Author: b-shi <brianshi@amd.com> * Free swap/localwritebase sgprs before post-loop * Defer sgpr allocation to remove holds in sgpr pool. Add Origami library logic files for Cijk_Alik_Bljk_F4BS_MXA32_MXB32 (base, nta4, ntb4 variants). * Remove uneeded alignment and comment * Add more epilogue tests * Remove older origami library for f4bs * Reorder post-loop code blocks to after persistant loop Misc fixes * Fix build issues, relax longjump sgpr requirements * Fix GSU0 branch logic --------- --- 3f034bf: Add F4HS and F4SS Origami library logic for FP4→F16 and FP4→F32 GEMM (ROCm#35) --- Author: Majedul Sujon <85503863+msujon-AMD@users.noreply.github.com> * Add F4HS and F4SS Origami library logic for FP4→F16 and FP4→F32 GEMM - Add 6 new yaml files (F4HS, F4SS) across Origami, Origami_nta4, Origami_ntb4 - Update F4BS yaml files: AssertSummationElementMultiple 32→256 for K%256 enforcement - Add ("F4", "F4", "H", "S") to _validGEMMTypes and _HPATypes in Problem.py * Add F4HS test cases to subtile_mxfp4.yaml Add two new benchmark problem blocks for FP4→F16 (F4HS): - No-bias block: same wavetile and problem size coverage as F4SS - Bias epilogue block: BiasDataTypeList [s, h], relu/none activations * Add F4HS (FP4->Half) type support to Tensile client Add TypedGemm_F4_H_S typedef and corresponding reference CPU solver case so F4HS (FP4 input, Float16 output, Float compute) problems can be validated by the benchmark client. --------- --- d0bc8fd: Rewrite subtile-based scheduler. Fix DU>64 & enable very large MT (ROCm#36) --- Author: sebvince <115461989+sebvince@users.noreply.github.com> * Initial support for DU>256 * Renaming * add option to do DU=512 in the tests * blocked K-major for scale * Change scaleSet swap logic * Update print functions * Put scales after values for avoid race conditions * Fix tests * more test * tweak printschedule display * Add PGR2 in the yaml tests * Add new scaleGROp * comment out failing tests * Revert "comment out failing tests" This reverts commit 1f5802c. * Draft new logical scheduler * Refactoring * Add more test on step1 * Add more tests on step1 * add bf16 320x320 test * reduce step1 code * Simplify step1 logic * validate some step1 test * Fix partition 2x2 test * more step1 test * 320x320 BF16 test * Add test DU512 + partition2x2 * Simplify step1 code * Add step2 tests * Fix multi-partition step2 * Add step2 du512, 2x2 partition test * Use common algo for all numPartitions * Draft for step3 tests * remove useless tests * New GR algo (draft) * [Step3] Add more test * Iteration on GR * Display ordered GR list with granularities * More test * Add some comments * Disable by default debug logs * Getting rid of step naming * Start remove AnnotatedOp (still there in group pass) * Split dependency Ops * Add todo on place_GRs pass * Valid test_annotate_deps_1x1_partition_DU256 * Test output looking better (still WIP) * single dep for LR tooo * Add remove_cross_deps pass * Fix bugs in dependency pass * insert_gr_lr_inc pass * Add group_lr_gr pass * Add emit pass * Quick port of instruction Emit code * Move emit function to separate file * Refactoring instructionEmitter * Port vgprTile tracking * Reworking second pass (WIP) * Display unrolling requirement * Unrolling check on 2nd pass * Generic validation for assign_vgpr pass * Fix unroll * Add inst schedule in standalone mode * Use lrGran for vgprTile size calculation * Fix bug in emit pass (missing depencency) * PreMFMA path + non-duplication scale load * missing globalReadLDSBufferSwap for GR_INC scales * add wairlr_sync on all LR->GR dep * add waitgr_sync op * remove_unnecessary_gr_deps * Change LR dispatch algo a bit to avoid too many waitgr_sync * Avoid duplicated loads in emitter * Fix bug on gr_emit code * GrInc pass. fix duplicated insertion for B * Fix missing LR_inc for SA/SB * preloop, NLL, NGLL * Simplify preloop * minor changes * Move unroll logic to scheduler * minor changes * Fix unroll id bug on NLL / NGLL * Disable post GRINC for now * Remove commented code * Handle 1x4, 4x1 gr read gran * Fix vmcnt computation * Use correct grCount mapping * Revert in emit logic on buffer_load for PGR0 needs * Add bf16 version in standalone test * Fix LR_Inc insertion on DU>64 * Add subIterK/Partition comment to codegen * Fix issue in GrInc placement * Remove last_mt * Fix LR MT index bug with muli-partition * Disable early LDS size check when subtileImpl is on * Add pass to remove redundant LR deps + fixed issue on dependency annotation pass * Remove more LR redundant deps * Only insert wait_lr_sync on deps * Simple algo to select partition config * Remove HC value for partitions... * Take into account all inflight GR (all tensors) * Fix tests and regressions on gr counts * Fix grCount merge calculation * Better display of dependencies * Add remove_wait_lr_sync after grouping * Add temporary non reg file * Change merge logic on GR grouping pass * Fix non necessary wait_lr_sync * Downgrade some waitlr_sync to sync + added 384x256 no reg test * non reg test 320x320 * Add larger MT * non reg test for fp4 256x256 * Moving out instructionScheduler * Remove old scheduler * Renaming scheduler * Re-work test * Add larger MT test cases * Rename non-ref test * Re-add standalone mode * Refactor DepOp * Remove dead code * Remove MFMATileSize class * Remove from_til_info * Avoid redundant tensor list creation * Remove hardcode granularities in vgrpTile allocation pass. Simplify code. * Re-enable # PGR=2 WG 4x1/1x4, K > DU tests * Remove unused GRScaleOp * DepRef renaming * Get rid of MT string representation * Remove TODO * EmmitedModule simplication * Use explicit pass dependencies * Renaming LogicalScheduler * Remove old test_InterleavingScheduler.py file * Commenting failing test for now * Remove debug logs * Disable lds padding when using UseSubtileImpl --- e8e8c09: Fix LR-GR dependency issue when DU>64 (ROCm#40) --- Author: sebvince <115461989+sebvince@users.noreply.github.com> * Fix and simplify logic for remove_unnecessary_lr_deps * Add new ref tests for 128x128x(128,64) --- 4aa441a: Rebase fix --- Author: b-shi <brianshi@amd.com> --- 5ba911e: Skip rocRoller for FP4-A/FP4-B + pre-swizzled scale layout (ROCm#42) --- Author: Archana Ramalingam <98564406+archana-ramalingam@users.noreply.github.com> --- 9c74998: Rebase fix --- Author: b-shi <brianshi@amd.com> --- 842b149: Addressed review comments for KernelWriter and KernelWriterAssembly --- Author: Koji Nakajima <knakajim@amd.com> --- dce43b1: Fix computeLoadSrd issue --- Author: Brad Nemanich <Brad.Nemanich@amd.com> --- c075bbf: Fix preSolution CPU re-sync regressing subtile_mxfp4.yaml --- Author: Brad Nemanich <Brad.Nemanich@amd.com> --- ced840f: Fix computeLoadSrd issue (ROCm#43) --- Author: bnemanich <brad.nemanich@amd.com> --- bc2f6dd: Small update for gfx950 mx tests + more - enable UseSubtileImpl for all gfx950 non subtile mx tests - skip all gfx950 mxfp8 - use MXScaleFormat=1 as default - set AssertSummationElementMultiple=256 for subtile mxfp4 - fix isSwizzledSubtile in computeLoadSrd --- Author: Koji Nakajima <knakajim@amd.com> --- 5c794b7: Fix gsuasb.yaml failures --- Author: b-shi <brianshi@amd.com> --- 727f8db: tensilelite: add solution reject conditions for UseSubtileImpl=1 (ROCm#38) --- Author: Majedul Sujon <85503863+msujon-AMD@users.noreply.github.com> --- 8928fbb: Add more reject conditions for Subtile --- Author: Koji Nakajima <knakajim@amd.com> --- 6e63ab6: Fix kringshift test failures --- Author: b-shi <brianshi@amd.com> --- b3e9724: Update reject condtion for DepthU in subtile case. Plus, update DepthU setting for gfx950 mx test cases --- Author: Koji Nakajima <knakajim@amd.com> --- 5ab6009: Fix build errors --- Author: Brad Nemanich <Brad.Nemanich@amd.com> --- 4a4edca: Update more mxfp4 tensilelite test cases --- Author: Koji Nakajima <knakajim@amd.com> --- bbbc553: Update change log --- Author: Brad Nemanich <Brad.Nemanich@amd.com> --- 6476c04: Add more reject conditions for gfx950 subtile --- Author: Koji Nakajima <knakajim@amd.com> --- c5828c4: Updated gfx950 mxfp4 test cases - add StreamK setting - skip groupgemm tests for now (groupgemm does not support streamK) --- Author: Koji Nakajima <knakajim@amd.com> --- f1fc2f1: Fix hipblaslt build error of gfx950 --- Author: Koji Nakajima <knakajim@amd.com> --- 70cea1b: Updated subtile_mxfp4.yaml (add StreamK) --- Author: Koji Nakajima <knakajim@amd.com> --- c1c9b2a: Add uninit lsc,lsp, etc.. fields for subtile --- Author: b-shi <brianshi@amd.com> --- c0c1f72: Fixed merge error in testing_matmul.hpp --- Author: Koji Nakajima <knakajim@amd.com> --- 191e0cb: Add missed batch_count >1 changes --- Author: archana-ramalingam <Archana.Ramalingam@amd.com> --- 01c52f8: Addressed PR comments --- Author: Koji Nakajima <knakajim@amd.com> --- 4e89c91: Reduce mxfp4 test time --- Author: Brad Nemanich <Brad.Nemanich@amd.com> --- 3dac20f: Prevent overflow for wgmxcc sgpr allocation --- Author: b-shi <brianshi@amd.com> --- 18dec79: Fix error with problem type --- Author: Brad Nemanich <Brad.Nemanich@amd.com> --- 9e69ffd: Add a reject conditoin for gfx950 mx + non Subtile --- Author: Koji Nakajima <knakajim@amd.com> --- 0eed3ba: Add more valid GEMM types --- Author: Brad Nemanich <brad.nemanich@amd.com> --- 8b5514e: Fix missing b build error --- Author: archana-ramalingam <Archana.Ramalingam@amd.com> --- f981ff5: Fix 1250 tests --- Author: Brad Nemanich <brad.nemanich@amd.com> --- d1e69d9: Add more FP4 tests --- Author: Brad Nemanich <Brad.Nemanich@amd.com> --- e3a688f: Add MXScaleFormat: 1 to all gfx950 mx test yaml --- Author: Koji Nakajima <knakajim@amd.com> --- aaef3f5: Add DataTypeMXSA,B setting in gfx950 mxfp4 logic yaml --- Author: Koji Nakajima <knakajim@amd.com> --- 861ef8e: Add DataTypeMXSA,B setting in gfx950 mxfp4 logic yaml (nta4,ntb4) --- Author: Koji Nakajima <knakajim@amd.com> Co-authored-by: Archana Ramalingam <Archana.Ramalingam@amd.com> Co-authored-by: Brad Nemanich <Brad.Nemanich@amd.com> Co-authored-by: Brian Shi <Brian.Shi@amd.com> Co-authored-by: James Newling <James.Newling@amd.com> Co-authored-by: Koji Nakajima <Koji.Nakajima@amd.com> Co-authored-by: Majedul Sujon <Majed.Sujon@amd.com> Co-authored-by: Sebastien Vince <Sebastien.Vince@amd.com> Co-authored-by: T.J. Alumbaugh <T.J.Alumbaugh@amd.com> ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Archana Ramalingam <Archana.Ramalingam@amd.com> Co-authored-by: Brad Nemanich <Brad.Nemanich@amd.com> Co-authored-by: Brian Shi <Brian.Shi@amd.com> Co-authored-by: James Newling <James.Newling@amd.com> Co-authored-by: Koji Nakajima <Koji.Nakajima@amd.com> Co-authored-by: Majedul Sujon <Majed.Sujon@amd.com> Co-authored-by: Sebastien Vince <Sebastien.Vince@amd.com> Co-authored-by: T.J. Alumbaugh <T.J.Alumbaugh@amd.com>
Saiyang-Zhang
pushed a commit
that referenced
this pull request
May 26, 2026
## Motivation https://amd-hub.atlassian.net/browse/AIHPBLAS-1467 ## Technical Details Fixed multiple issues preventing TensileLight from correctly generating and executing kernels when UseBeta=false (beta parameter not used in GEMM operations). Enabled bounds checking validation to work correctly with this configuration. Files Modified 1. Tensile/KernelWriterAssembly.py Issue: KeyError when accessing Beta SGPR register when UseBeta=false Fix: Added conditional check before accessing Beta SGPR if kernel["ProblemType"]["UseBeta"]: moduleExternalArgs.addComment("Read Beta") moduleExternalArgs.addModuleAsFlatItems(self.externalArgLoader.loadAllKernArg( self.sgprs["Beta"], "KernArgAddress", self.states.numSgprBeta)) 2. Tensile/SolutionStructs/Problem.py Issue: UseBeta serialized as integer (0/1) instead of boolean in YAML, causing C++ parser errors Fix: Ensure UseBeta is always stored as boolean self.state["UseBeta"] = bool(self.state["UseBeta"]) 3. client/src/ReferenceValidator.cpp Issue #1: Buffer allocation check didn't verify if buffer pointer was valid Fix: Check both size and pointer validity // Only skip reallocation if size matches AND buffer is valid if(m_cpuResultBufferSize == bytes && m_cpuResultBuffer.get() != nullptr) return; Issue ROCm#2: hipFree compiler warning about nodiscard attribute Fix: Cast return value to void in lambda deleter uint8_t* buffer; HIP_CHECK_EXC(hipHostMalloc((void**)&buffer, bytes, 0)); m_cpuResultBuffer.reset(buffer, [](uint8_t* p) { (void)hipFree(p); }); Issue ROCm#3: Attempting to validate null/empty tensors Fix: Skip validation for null pointers or zero-sized tensors // Skip validation if pointers are null or maxElements is 0 if(resPtr == nullptr || refPtr == nullptr || result.maxElements[i] == 0) { if(Debug::Instance().printTensorInfo()) std::cout << "Skipping validation for tensor " << tensor.getName() << std::endl; continue; } Issue ROCm#4: Trying to copy padding bytes from output tensors that don't have padding Fix: Only use maxElement for input tensors // For output tensors, don't use maxElement with padding if(boundsCheck == BoundsCheckMode::NaN && !tensor.isOutput()) elementsToCopy = maxElement; Issue ROCm#5: Bounds checking validation on output tensors without padding buffers Fix: Skip bounds checking for output tensors // Only check bounds for input tensors (output tensors don't have padding buffers) if(boundsCheck == BoundsCheckMode::NaN && !tensor.isOutput()) 4. client/src/DataInitialization.cpp Issue: hipMemcpy with null pointers causing runtime errors Fix: Added null pointer check void* copyInputBuffers(const TensorDescriptor& descriptor, void* dst, void* src, size_t totalElements, hipMemcpyKind kind) { // Skip copy if no elements to copy or if pointers are null if(totalElements > 0 && dst != nullptr && src != nullptr) { HIP_CHECK_EXC(hipMemcpy(dst, src, descriptor.elementBytes() * totalElements, kind)); } return dst; } ROCm@0d6cd23 ## Test Plan NA ## Test Result ``` ========================================================================================== 105 passed, 83 skipped, 1 warning in 1040.87s (0:17:20) ========================================================================================== ``` ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Signed-off-by: pdhirajkumarprasad <dhirajp@amd.com>
EwanC
pushed a commit
that referenced
this pull request
Jul 8, 2026
## What
Make origami's installed `CTestTestfile.cmake` use paths relative to the
test
directory instead of `${CMAKE_CURRENT_LIST_DIR}`.
## Why
`${CMAKE_CURRENT_LIST_DIR}` is **empty** when `ctest` parses an
*installed*
`CTestTestfile.cmake` — its restricted parser does not populate that
variable
(unlike a normal configure or `cmake -P` run). The relocatable file
emitted by
ROCm#9047 used it for the test command, `WORKING_DIRECTORY`, and the
`LD_LIBRARY_PATH`/`PYTHONPATH` environment, so `WORKING_DIRECTORY`
expanded to
`""` and every test failed to launch:
```
Failed to change working directory to : No such file or directory
1/2 Test #1: origami-tests ....................***Not Run
2/2 Test ROCm#2: origami_python_tests .............***Not Run
```
ctest already runs each test with its working directory defaulting to
the
directory containing the `CTestTestfile.cmake`, including under
`ctest --test-dir <dir>` invoked from an unrelated cwd. So relative
paths
resolve correctly and relocatably without the (empty) variable. This
also aligns
origami with the sibling hipBLASLt installed test file, which uses the
same
relative-path convention ("Tests are defined with relative paths to work
in the
installed location").
## Change
`shared/origami/tests/CMakeLists.txt` — in the installed
`_ctest_content`:
- `"${CMAKE_CURRENT_LIST_DIR}/../origami-tests"` → `"../origami-tests"`
- pytest target `"${CMAKE_CURRENT_LIST_DIR}/tests"` → `"tests"`
- dropped `WORKING_DIRECTORY "${CMAKE_CURRENT_LIST_DIR}"` (default is
already the file's dir)
- env paths `${CMAKE_CURRENT_LIST_DIR}/../../lib[...]` →
`../../lib[...]`
## Verification
Reproduced and validated in
`ghcr.io/rocm/therock_build_manylinux_x86_64`
(ctest 3.27.9, python 3.12) against a real origami build. Both installed
`CTestTestfile.cmake` variants were generated by actual `cmake`
CONFIGURE from
the committed source and run via `ctest --test-dir` from an unrelated
cwd:
| | C++ `origami-tests` | `origami_python_tests` |
|---|---|---|
| Before (`${CMAKE_CURRENT_LIST_DIR}`) | Not Run | Not Run |
| After (relative) | Passed | Runs |
🤖 Generated with [Claude Code](https://claude.com/claude-code)
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
BalintCsala
pushed a commit
that referenced
this pull request
Jul 13, 2026
… uchar type (ROCm#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 ROCm#2: rpp_sanity_test_brightness_hip_f32 .......... Passed 7.75 sec Start 3: rpp_qa_tests_tensor_image_host_all 3/10 Test ROCm#3: rpp_qa_tests_tensor_image_host_all .......... Passed 42.43 sec Start 4: rpp_qa_tests_tensor_voxel_host_all 4/10 Test ROCm#4: rpp_qa_tests_tensor_voxel_host_all .......... Passed 2.26 sec Start 5: rpp_qa_tests_tensor_misc_host_all 5/10 Test ROCm#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 ROCm#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 ROCm#7: rpp_qa_tests_tensor_image_hip_all ........... Passed 84.44 sec Start 8: rpp_qa_tests_tensor_voxel_hip_all 8/10 Test ROCm#8: rpp_qa_tests_tensor_voxel_hip_all ........... Passed 4.62 sec Start 9: rpp_qa_tests_tensor_misc_hip_all 9/10 Test ROCm#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 ROCm#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.
BalintCsala
pushed a commit
that referenced
this pull request
Jul 23, 2026
…hipcxx (#1) * Remove duplicates from inc/gpu/__algorithm * Remove inc/gpu/__atomic * Remove __memory duplicates * Update __memory usages * Remove __iterator duplicates * Remove __numeric duplicates * Remove __type_traits duplicates * Remove __functional duplicates * Update __functional usage in __thread/worknode.h * Remove cstring duplicate * Remove cstdlib duplicate * Remove __tuple_dir duplicate * Remove __utility duplicates * Update __utility usage * Remove __split_buffer header * Remove vector header * Remove memcmp.h from __clib * Remove __string dir * Remove device unique_ptr tests * Remove host unique_ptr tests * Remove exposed unique_ptr_test.cxx * Remove vector_test * Update compile_test.cxx * Remove atomic_test.cxx * Remove atomic lit tests --------- Co-authored-by: Marko Savic <marsavic@amd.com> [ROCm/hipthreads commit: 6e7b1eb]
BalintCsala
pushed a commit
that referenced
this pull request
Jul 23, 2026
…hipcxx (#1) * Remove duplicates from inc/gpu/__algorithm * Remove inc/gpu/__atomic * Remove __memory duplicates * Update __memory usages * Remove __iterator duplicates * Remove __numeric duplicates * Remove __type_traits duplicates * Remove __functional duplicates * Update __functional usage in __thread/worknode.h * Remove cstring duplicate * Remove cstdlib duplicate * Remove __tuple_dir duplicate * Remove __utility duplicates * Update __utility usage * Remove __split_buffer header * Remove vector header * Remove memcmp.h from __clib * Remove __string dir * Remove device unique_ptr tests * Remove host unique_ptr tests * Remove exposed unique_ptr_test.cxx * Remove vector_test * Update compile_test.cxx * Remove atomic_test.cxx * Remove atomic lit tests --------- Co-authored-by: Marko Savic <marsavic@amd.com>
BalintCsala
pushed a commit
that referenced
this pull request
Aug 7, 2026
…eaders (ROCm#10425) ## Summary Nine `switch` statements in `projects/rocalution/src/base/host/host_io.cpp` pick the conversion to apply when a matrix file's value type differs from the in-memory `ValueType`. In all nine, the `complex32` case has a body but no `break`, so control falls into the `complex64` case: ```cpp case rocsparseio_type_complex32: { copy_mixed_arrays(nnz, val[0], (const std::complex<float>*)tmp_val); } case rocsparseio_type_complex64: { copy_mixed_arrays(nnz, val[0], (const std::complex<double>*)tmp_val); break; } ``` The result is a heap buffer over-read plus silent corruption of the values that were just converted correctly. Related: found during the same monorepo sweep as ROCm#10413. ## Why it is a bug Three facts combine, all in-tree: 1. **The scratch buffer is sized from the file's type, not the in-memory type.** `host_io.cpp:918-922`: ```cpp size_t sizeof_val_type; status = rocsparseio_type_get_size(file_val_type, &sizeof_val_type); tmp_val = malloc(nnz * sizeof_val_type); ``` and the `switch` is on `file_val_type`, so reaching the `complex32` case means the buffer is `nnz * 8` bytes. 2. **The two widths differ by 2x.** `src/utils/rocsparseio.hpp:415-421`: ```cpp case complex32: return sizeof(float) * 2; // 8 case complex64: return sizeof(double) * 2; // 16 ``` 3. **`copy_mixed_arrays` reads `size` elements of its source type.** `host_io.cpp:656-667`: ```cpp for(size_t i = 0; i < size; ++i) { x[i] = static_cast<X>(y[i]); } ``` So the fallthrough reads `nnz * 16` bytes from an `nnz * 8` byte allocation: an `nnz * 8` byte over-read. It then writes those garbage values over `val[0]`, discarding the correct conversion the `complex32` case had already performed. ## Reachability, stated precisely Exactly one configuration reaches the fallthrough: **a file whose value type is `complex32`, read into `ValueType` `std::complex<double>`.** I checked the other possibilities rather than assuming: | `ValueType` | What happens | Reaches fallthrough | | --- | --- | --- | | `std::complex<float>` | `required_val_type == complex32 == file_val_type`, so `same_val_type` is true and the whole `switch` is skipped | no | | `std::complex<double>` | `complex32` case runs the real `complex<double> <- complex<float>` conversion, then falls through | **yes** | | `int8_t`, `float`, `double` | the `complex32` case hits a `copy_mixed_arrays` specialization whose body is `throw 1;`, so it throws first | no | The throwing specializations are at `host_io.cpp:696-731`. `std::complex<double>` is explicitly instantiated for every one of these readers (`host_io.cpp:4241`, `4303`, `4361`, `4434`, `4504`, and the CSR/MCSR/BCSR equivalents), so the path is compiled into the shipping library. It is the ordinary mixed-precision case of loading a single-precision complex matrix into a double-precision complex one. I want to be straight that this is one specific type pairing rather than "any mismatched complex read". It is still a live memory-safety defect on a supported and instantiated path. ## Scope Affected readers, one site each except `hyb` which has two (the COO value array and the ELL value array): ``` host_io.cpp:1022 read_matrix_csr_rocsparseio switch on file_val_type host_io.cpp:1330 read_matrix_mcsr_rocsparseio switch on file_val_type host_io.cpp:1681 read_matrix_bcsr_rocsparseio switch on file_val_type host_io.cpp:1967 read_matrix_coo_rocsparseio switch on file_val_type host_io.cpp:2231 read_matrix_dia_rocsparseio switch on file_val_type host_io.cpp:2488 read_matrix_ell_rocsparseio switch on file_val_type host_io.cpp:2858 read_matrix_hyb_rocsparseio switch on file_coo_val_type host_io.cpp:2931 read_matrix_hyb_rocsparseio switch on file_ell_val_type host_io.cpp:3123 read_matrix_dense_rocsparseio switch on file_val_type ``` All nine switches have identical shape. Every case group ends in `break` except `complex32`: ``` switch@996 on file_val_type: break case int32+int64 break case int8 break case float32 break case float64 NO-EXIT case complex32 break case complex64 ``` **`complex32` also appears 13 more times as a stacked case label sharing a body with its neighbours**, for example `host_io.cpp:955-962`: ```cpp case rocsparseio_type_int8: case rocsparseio_type_float32: case rocsparseio_type_float64: case rocsparseio_type_complex32: case rocsparseio_type_complex64: { break; } ``` Those are correct and are **not** touched. Only the nine cases that carry their own body are changed. A first pass of my own scan counted those 13 as defects before I distinguished stacked labels from bodies, so the number in this PR is 9, not 22. ## Fix `break;` added to the nine `complex32` bodies. Nine lines, no other change: ```diff case rocsparseio_type_complex32: { copy_mixed_arrays(nnz, val[0], (const std::complex<float>*)tmp_val); + break; } ``` ## Verification No AMD GPU was used, and none is needed: this is host-side matrix-file I/O. I built a reduced standalone reproducer modelling the buffer sizing, the 8 versus 16 byte type widths, the `copy_mixed_arrays` overload set including the throwing specializations, and the `switch` itself, on the reachable `ValueType = std::complex<double>` instantiation. `nnz = 4`, so `tmp_val` is 32 bytes and the fallthrough reads 64. **Before, `clang++ -std=c++17 -g -O0 -fsanitize=address`:** ``` tmp_val = malloc(4 * 8) = 32 bytes after complex32 case: val[0] = (1,10) (2,11) (3,12) (4,13) <-- correct ================================================================= ==32738==ERROR: AddressSanitizer: heap-buffer-overflow on address 0x603000001c50 at pc 0x000103082c84 bp 0x00016d5d6120 sp 0x00016d5d58d0 READ of size 16 at 0x603000001c50 thread T0 #0 0x000103082c80 in __asan_memcpy+0x400 (libclang_rt.asan_osx_dynamic.dylib:arm64e+0x3ec80) #1 0x0001028290f0 in void copy_mixed_arrays<std::__1::complex<double>, std::__1::complex<double>>(unsigned long, std::__1::complex<double>*, std::__1::complex<double> const*) repro.cpp:53 ROCm#2 0x000102828b58 in main repro.cpp:129 ROCm#3 0x00018c759d50 in start+0x1c0c (dyld:arm64e+0x8d50) 0x603000001c50 is located 0 bytes after 32-byte region [0x603000001c30,0x603000001c50) allocated by thread T0 here: #0 0x000103085164 in malloc+0x78 (libclang_rt.asan_osx_dynamic.dylib:arm64e+0x41164) #1 0x000102828980 in main repro.cpp:82 ROCm#2 0x00018c759d50 in start+0x1c0c (dyld:arm64e+0x8d50) ``` Note the frame: the over-read is attributed to the `copy_mixed_arrays<complex<double>, complex<double>>` instantiation, which is the generic template reached only through the fallthrough. At `-O1` the loop vectorises and ASan reports the same overflow as a single `READ of size 64`. **Before, no sanitizer**, showing the second half of the defect: ``` tmp_val = malloc(4 * 8) = 32 bytes after complex32 case: val[0] = (1,10) (2,11) (3,12) (4,13) <-- correct after complex64 case: val[0] = (524288,1.04858e+06) (2.09715e+06,4.19431e+06) (0,0) (0,0) <-- clobbered done ``` The correct values are replaced by the bit patterns of pairs of `std::complex<float>` reinterpreted as single `std::complex<double>` values, trailed by the out-of-bounds tail. **After, `-DFIXED`, ASan clean at both `-O0` and `-O1`:** ``` tmp_val = malloc(4 * 8) = 32 bytes after complex32 case: val[0] = (1,10) (2,11) (3,12) (4,13) <-- correct done EXIT=0 ``` ### Why no compiler caught it `-Wall -Wextra` are silent. `-Wimplicit-fallthrough` does flag it, but clang does not imply it from either, and rocALUTION's CMake never enables it. Turning it on for `projects/rocalution` would prevent recurrence, and I am happy to do that as a follow-up if you want it separated from this fix. ## A note on the unit test policy check `tools/libraries_pr_bot/policy.yml` requires an accompanying test file for `.cpp` changes, with `exempt_paths: []`. I have not added one, and I would rather explain why than route around the check. `projects/rocalution/clients/tests/` has 31 gtest files and **none of them cover file I/O at all**: `git grep -l "ReadFileRSIO\|WriteFileRSIO\|rocsparseio" -- projects/rocalution/clients/` returns nothing. So there is no existing test file to extend, and a new one needs a rocsparseio fixture written at a complex32 value type, a new entry in `clients/tests/CMakeLists.txt`, and `init_rocalution()`. I cannot build or run the rocALUTION client suite here, and shipping unexercised CMake and gtest code alongside a memory-safety fix would make this PR harder to trust, not easier. Two things I am happy to add, whichever you prefer: 1. A `WriteFileRSIO` then `ReadFileRSIO` round-trip test in a new `clients/tests/test_host_io.cpp`, writing a small `complex<float>` matrix and reading it back into a `LocalMatrix<std::complex<double>>`, asserting the values match. That is the faithful regression test and it fails on `develop` under ASan. I would need someone to run it once, since I have no ROCm runtime. 2. Enabling `-Wimplicit-fallthrough` for `projects/rocalution`, which pins the whole class rather than this instance. I can also hand over the standalone reproducer if it is useful for the review. ## Checklist - Based on `develop` (`86a82b0`), single commit, DCO signed off. - Re-verified against current `develop` HEAD rather than an earlier snapshot: 9 affected sites, 13 stacked-label sites correctly excluded, all other cases in all nine switches confirmed to still end in `break`. - Prior art checked: no open PR or issue mentions `host_io`, `copy_mixed_arrays` or `rocsparseio_type_complex32`, and `git log -S 'rocsparseio_type_complex32'` shows no partial fix. - `projects/rocalution` is **not** in the root `.pre-commit-config.yaml` exclude list, so `clang-format` 18.1.4, the pinned version, was run against `projects/rocalution/.clang-format`. It reformats nothing beyond the nine added lines. No trailing whitespace, file still ends in a newline. - Behaviour is unchanged for every configuration except the one that was reading out of bounds. Signed-off-by: Aditya Singh <adisin650@gmail.com>
BalintCsala
pushed a commit
that referenced
this pull request
Aug 26, 2026
ISSUE ID: ROCm#8997 ## Motivation The TileEngine → Dispatcher bridge had no path for the weight-preshuffled GEMM op **gemm_preshuffle**, which pre-permutes the B (weight) tensor into the pipeline's packed layout for higher throughput on weight-heavy GEMMs. This is a real Old-TE capability with no dispatcher equivalent, so this PR adds the bridge so the dispatcher can generate and launch preshuffle GEMM at parity with the legacy Tile Engine version. The capability set matches the Old-TE op exactly: `fp16`, `bf16`, `fp8`, `bf8`, `rcr` layout only (A row-major, B col-major, C row-major), with the B tensor preshuffled into the packed layout. It follows the same bridge scheme as ROCm#8997 (regular GEMM), ROCm#9000 (grouped), and ROCm#9028 (stream-K). ## Test Plan - Run the CPU-only unit tests (no GPU required): `python3 -m pytest dispatcher/tests/test_preshuffle_bridge.py -v` - Codegen + build the `.so` for all 4 dtypes and verify output vs an fp32 NumPy reference at M=N=K=512 on gfx942 / MI300X. - Confirm the preshuffle B-shuffle matches `ck_tile::shuffle_b` byte-for-byte. ## Test Result - CPU-only unit tests pass (8 passed). - Codegen + `.so` build succeed for all 4 dtypes (0 codegen failures). Output verified vs fp32 reference at M=N=K=512: | dtype | max_rel | status | |-------|---------|--------| | fp16 | 5.2e-4 | PASS | | bf16 | 5.5e-3 | PASS | | fp8 | 3.4e-2 | PASS | | bf8 | 6.5e-2 | PASS | - The preshuffle B-shuffle (the #1 correctness risk) is verified byte-for-byte via `ck_tile::shuffle_b`; the permuteN shuffle gives ~1.25 rel error and is deliberately not used, since the bridge does not emit the permuteN pipeline. > **Note (arch support):** the correctness above is on gfx942 (MI300X). fp8/bf8 preshuffle is unsupported on **gfx950** for *both* this bridge and Old-TE — the underlying kernel fails to launch there (`status -2`) — while fp16/bf16 are correct on both arches (see the cross-arch note in the comments). --- **Related PRs / references (TileEngine → Dispatcher GEMM bridge series):** ROCm#8997 (regular GEMM fp16/bf16 all-layout), ROCm#9000 (grouped GEMM), ROCm#9028 (stream-K), ROCm#8887 (fp8/bf8/int8). This PR is a sibling in the same bridge effort tracked across those PRs. --------- Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> Co-authored-by: Thrupti Raj Lakshmana Gowda <thruptiraj.lakshmanagowda@amd.com>
BalintCsala
pushed a commit
that referenced
this pull request
Sep 16, 2026
…Cm#11978) ## Summary Completes the library carve for convolution, following the attention split (ROCm#8977). Every convolution kernel builder, dispatch policy, test, parity emitter, benchmark, and builder harness now lives under `rocke/library/`. `rocke/platform/` keeps none of them — not even re-export shims — except where a real architectural constraint (the manifest-runner's lazy-resolution registry) requires a thin pointer. Also ports the two-stage deterministic wgrad feature (ROCm#10571) into the library, since that PR landed on `develop` after this branch was cut, and fixes a handful of bugs found in review (see "Bug fixes" below). This is a full move, not a copy: `git diff --diff-filter=D` against the merge-base shows the platform originals gone, not left behind. An earlier version of this PR left the platform copies in place as unreferenced duplicates; that duplication has been removed in this revision. JIRA-ID: AICK-2047 --- ## What changed and why ### Kernel builders (`library/kernels/common/` and per-arch subpackages) | File | Description | |---|---| | `conv_implicit_gemm.py` | Forward NHWC implicit-GEMM | | `conv_implicit_gemm_wgrad.py` | Backward-weight implicit-GEMM; adds `two_stage` / `force_deterministic` fields + workspace-store epilogue (ported from ROCm#10571) | | `conv_implicit_gemm_wgrad_two_stage.py` | Two-stage pipeline launcher (ported from ROCm#10571) | | `conv_wgrad_workspace_reduce.py` | Stage 2 workspace-reduce kernel builder (ported from ROCm#10571) | | `conv_implicit_gemm_dgrad.py` | Backward-data implicit-GEMM | | `conv_direct_grouped.py` | Direct grouped conv (4c / 8c / 16c / 32c / depthwise variants) | | `deep_fused_conv_pool.py` + gfx950 / gfx1151 / gfx1201 arch variants | Fused conv+pool | | `img2col.py` | Im2col transform | All `build_*` functions take `arch` as a keyword-only parameter, per the builder contract established in ROCm#11237. ### Manifest runner `platform/python/rocke/run_manifest.py` no longer imports the conv manifest runner eagerly at module scope. It resolves `conv_fp16` / `conv_bf16` / `conv_fp32` lazily through the existing `_LIBRARY_RUNNER_MODULES` registry (the same pattern already used for `deep_fused_conv_pool_fp16` / `_i8i4`), pointing at `kernels.common.manifest_runner.conv`. This is the one place `platform` still names a `library` module path, and it does so only inside a function body, resolved at call time — `platform` stays standalone-installable with `library` absent. ### Tests | File | Description | |---|---| | `library/tests/test_conv_fwd_correctness.py` | GPU correctness for fwd | | `library/tests/test_conv_wgrad_correctness.py` | GPU correctness for wgrad; includes `TestConvWgradTwoStage` (29 methods, ported from ROCm#10571) | | `library/tests/test_conv_dgrad_correctness.py` | GPU correctness for dgrad | | `library/tests/test_direct_conv_correctness.py` | GPU correctness for direct-grouped conv (moved from `platform/tests/instances/`) | | `library/tests/test_conv_build_signatures.py` / `test_builder_signature_contract.py` | Build-signature contracts (no GPU) | | `library/tests/parity/conv_implicit_gemm_{,wgrad,dgrad}_emit.{py,c}` | Python/C byte-identity emitters; wgrad extended with the two-stage and split-K configs ported from ROCm#10571 | | `library/tests/parity/conv_wgrad_workspace_reduce_emit.{py,c}` | Parity configs for the Stage 2 reduce kernel | | `library/tests/parity/conv_direct_grouped_emit.{py,c}` | Extended to cover all config variants (16c/4c/8c/32c/depthwise) that platform's copy had and library's did not | | `library/tests/parity/img2col_emit.{py,c}` | Moved from `platform/tests/instances/parity/` — it already imported `kernels.common.*` at module scope and had no library-side counterpart | | `library/tests/dispatch/test_grouped_conv_wgrad_dispatch.py` | Wgrad dispatch selection | `library/tests/test_conv_dgrad_correctness.py` resolves `platform`/`library` roots via `rocke.assets.library_root()` / `platform_root()` instead of hand-rolled `parents[N]` math (the old math pointed one level too high and silently ran its subprocess benchmark with a broken `PYTHONPATH`). The child process's own `PYTHONPATH` is now built by prefixing the current process's `sys.path`, so the same test works unchanged whether it's running from a source checkout or an installed CTest artifact. ### Benchmarks / builders `benchmark_implicit_gemm_conv.py`, `benchmark_direct_conv.py`, `conv_reference.py`, `benchmark_conv_compare.py`, and the gfx950 / gfx1151 / gfx1201 deep-conv-fusion builder harnesses all live under `library/benchmarks/common/` and `library/builders/<arch>/` now. `benchmark_implicit_gemm_conv.py` includes the `--two-stage` sweep support (CLI flag, Stage 1 + Stage 2 GPU run, results reporting) that had been left out of an earlier draft of this move. `library/builders/common/conv_reference.py` now slices the weight tensor to `C // groups` channels before the torch reference conv, matching platform's original grouped-conv fix — without it, reference values for `groups > 1` were wrong. ### Documentation Fixed every dangling reference to the old `platform/python/rocke/instances/common/` and `platform/tests/instances/` conv paths across `dsl_docs/`, builder `README.md`s, and in-code docstrings (Run: instructions, "Key Files" tables, prose) — both broken markdown links and plain-text mentions. ### CI - `rocke_library_conv_pytest` (existing CTest target): CPU-only dispatch/signature tests, no GPU required. - `rocke_library_conv_gpu_pytest` (new): registers `test_conv_{fwd,wgrad,dgrad}_correctness.py` and `test_direct_conv_correctness.py`. These were collected by the platform-wide `rocke_pytest` target while they lived under `platform/tests/instances/`; moving them to `library/tests/` put them inside that target's `--ignore=tests/library` scope, silently dropping them from CI. They self-skip when no supported GPU/torch is detected, so a host-only runner still passes; the `gpu` label lets a runner select or exclude them explicitly. GPU-nightly runner registration (which machines actually run the `gpu`-labelled set) is out of scope here — no such runner config exists anywhere in this repo for me to safely edit; that needs a human with access to that infra. ### Bug fixes found during review and re-verification - `platform/python/rocke/core/lower_cktile.py` deferred its `kernels.common.conv_*` import until after the `UniversalGemmSpec` branch — it was unconditionally importing a `library`-layer module before checking which spec type it had, breaking plain-GEMM lowering on a `platform`-only install. - `platform/python/rocke/examples/common/hip_lowering_parity.py` had a module-level `kernels.common.*` import that broke standalone-platform installs for all ~40 non-conv families this audit tool also covers; deferred into the function body that actually needs it. - Both `test_direct_conv_correctness.py` and `test_conv_fwd_correctness.py` had a pre-existing HIP-context race (predates this migration, confirmed via `git blame` against ROCm#11168): calling rocke's HIP runtime before torch touches CUDA claims the process HIP context and permanently breaks `torch.cuda()` for the rest of the process. Fixed by having torch claim the context first when present. - Deleted a stale `platform/python/rocke/examples/gfx1151/deep_conv_fusion/` copy whose `compare_prebuilt.py` imported a `run_manifest` symbol that no longer exists (removed when the manifest runner was made lazy); the library copy was already the correct, current one. --- ## Verification - **Byte-identity gate** (the repo's #1 invariant): GREEN at both `llvm20` and `llvm22`, all families, after every kernel-file move. - **Platform pytest** (standalone, `library` absent from `sys.path`): 853 passed, 64 skipped, 1548 subtests passed. - **Library CPU tests** (layering, build-signature contract, dispatch): 326 passed, 216 subtests passed. - **GPU correctness, gfx942 (MI300X), run locally in this environment**: `test_conv_fwd_correctness.py`, `test_conv_wgrad_correctness.py`, `test_conv_dgrad_correctness.py`, `test_direct_conv_correctness.py` together — 73 passed, 68 skipped (arch-gated: gfx950/gfx1250-only variants not exercisable on this hardware), 366 subtests passed, ~5.5 minutes. - **Standalone-install invariant**: verified with a script that strips `library` from `sys.path` and imports every `rocke.*` submodule via `pkgutil.walk_packages` — zero failures. - **No architecture violations**: `grep -rn "rocke\.instances" library/` returns only pre-existing, unrelated attention/matmul_nbits references; zero `kernels`/`builders`/`dispatch` imports at module scope anywhere under `platform/python/`. ## Known, pre-existing gaps (not introduced or hidden by this PR) - **No dgrad dispatch policy layer.** Unlike fwd/wgrad, there is no `dispatch/conv_dgrad.py` request-routing module, in either tree, with any git history — dgrad is built directly today. Writing one needs real dispatch-policy design input; out of scope here rather than invented blind on a shipping path. - **gfx950 hardware was not available in this environment.** gfx950-specific conv paths (deep-fused-conv-pool, WMMA variants) are byte-identity-checked but not GPU-run here; they need a human with gfx950 access to sign off numerically. --------- Co-authored-by: Claude <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
20 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Merge latest ROCm/rocm-libraries develop branch to users/EwanC/BNFwdInf_InvVar to fix CI for ROCm#3619.
Resolved conflicts for: