[origami] Major refactoring of codebase by neoblizz · Pull Request #1859 · ROCm/rocm-libraries

neoblizz · 2025-09-27T03:37:12Z

Important

Looking for feedback on the refactored APIs and C++ best-practices.

Refactor

Python APIs & Unit Tests

Update python APIs + ATEN compatible APIs
Update the tests
Update rocroller @minsukim-amd
hipblaslt/tensilelite/scripts to use the new API @minsukim-amd

Testing Infrastructure

Replace YAML-based tests with Catch2 C++ tests
Replace GTest with Catch2 framework
Convert YAML-driven parameterized tests to pure C++ tests
Update common.hpp with direct C++ helper functions
Update CMakeLists.txt to use Catch2 instead of GTest
Remove Boost dependency (was only for YAML)
All test data now hardcoded in C++ for type safety
Better IDE support, debugging, and error messages

Questions

Rocroller: coordinate on the API design, should they be reused?
Reuse dim3_t from hip_runtime.h (does HIP's dim3 have the same functionality)?
Should transpose/data-types be part of config as well for precomputing?

Motivation

Make it simpler to make model updates.
Well-defined, scalable, clean interface.

Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

- New file , consolidates various origami types. - New structs to replace the growing tuple. - API updates to make it scalable, down from 20+ parameters to a few. - Lots of redundant code removal. '

ryanswann-amd · 2025-09-29T18:42:44Z

+    # Select reduction strategy  
+    grid_algorithm = origami.grid_selection_t.analytical  # default to analytical
+    if args.dynamic_grid_version == 0:
+        grid_algorithm = origami.grid_selection_t.number_of_cus
+    elif args.dynamic_grid_version == 1:
+        grid_algorithm = origami.grid_selection_t.min_resources
+    elif args.dynamic_grid_version == 2:
+        grid_algorithm = origami.grid_selection_t.energy_aware
+    elif args.dynamic_grid_version == 3:
+        grid_algorithm = origami.grid_selection_t.reduction_cost_aware
+    elif args.dynamic_grid_version == 4:
+        grid_algorithm = origami.grid_selection_t.data_parallel
+    elif args.dynamic_grid_version == 5:
+        grid_algorithm = origami.grid_selection_t.analytical
+    elif args.dynamic_grid_version == 6:
+        grid_algorithm = origami.grid_selection_t.k_split_aware
+


Is this an algorithm or a heuristic?

ryanswann-amd · 2025-09-29T18:50:21Z

+  /// Transpose types (TT, TN, NT, TT.)
+  bool transpose_a;
+  bool transpose_b;


I've always thought this nomenclature was kind of weird. Could we use something like "contiguous dimensions" or "dimensions contiguous in memory"? We use that information later on to do transpose specific stuff based on how it's aligned in the memory hierarchy

bnemanich · 2025-09-30T16:35:15Z

+  data_type_t d_dtype;
+
+  /// Compute type.
+  data_type_t mi_dtype;


Is this the same as the accumulation type? On gfx12, there are multiple accumulation types that can be used. We might want to add that to future proof in case we start supporting gfx12.

bnemanich · 2025-09-30T16:38:41Z

+  std::size_t batch;
+
+  /// Transpose types (TT, TN, NT, TT.)
+  bool a_transpose;


Might make sense to use an enumeration here. This is how hipBLASLt supports swizzling the inputs, which we might want to support.

@bnemanich as in an enum: TT, TN, NT, NN?

bnemanich · 2025-09-30T19:13:05Z

-            }
-        };
+struct matrix_instruction {
+  size_t MI_M;


Using all capitol letters doesn't match with formatting. Also, shouldn't need MI in the names, since they are part of matrix_instruction. Maybe we can just call the values m, n, and k.

bnemanich · 2025-09-30T19:19:13Z

+ public:
+  enum class architecture_t { gfx90a, gfx942, gfx950, Count };
+
+  static architecture_t arch_name_to_enum(const std::string& str) {


I think we should try to move source code out of the header files and move to source files when possible. Most of the code in this file can be moved to a source file.

This is why it was moved out of source file: #1862. I need to revisit this to make sure we don't cause this problem again.

Yes, I think that was due to the static values that were declared in this class. I actually think we should re-think a lot of this. Why are we using static methods everywhere in this file? Do we really need to have static variables as well? Those can cause weird issues, like the one pointed out in PR 1862.

pghysels · 2025-10-01T19:17:33Z

+
+  std::size_t hash() const {
+    return std::hash<size_t>()(MI_M) ^ std::hash<size_t>()(MI_N) ^ std::hash<size_t>()(MI_K) ^
+           std::hash<data_type_t>()(mi_input_type);


This creates some conflicts, since a hash of a size_t is usually the identity operation.
For instance I think matrix_instruction(32, 32, 4, data_type_t::Half) has the same hash as matrix_instruction(16, 16, 4, data_type_t::Half) and matrix_instruction(4, 4, 4, data_type_t::Half)

pghysels · 2025-10-01T19:19:11Z

+    // Optionally ensure we do not exceed LDS capacity
+    // (If you want to factor in WGM, add it to your check_lds_capacity signature.)
+    // For now, let's just check the tile itself:
+    if (!check_lds_capacity(hardware, mt, problem.a_dtype, problem.b_dtype)) {


I would move this out of the loop for now. Can always put it back when adding workgroupmapping as a parameter.

pghysels · 2025-10-01T19:57:59Z

-
-        return std::make_tuple(numActiveCUs, numWaves, splitFactor);
+    l2_A_uncached_reads = static_cast<long long>(l2_m) * config.mt.mk();
+    l2_B_uncached_reads = static_cast<long long>(l2_n) * config.mt.nk();


This doesn't need to be in the loop?

pghysels · 2025-10-01T19:59:09Z

+  }
+
+  // Total reads considering repeated usage
+  long long l2_A_reads = static_cast<long long>(l2_m) * l2_n * config.mt.mk();


This is long long l2_A_reads = l2_n * l2_A_uncached_reads; ?

bnemanich · 2025-10-02T14:17:42Z

+    std::cout << "===========================================================\n";
+  }
+  // Debug tracking info
+  mutable std::unordered_map<std::string, std::string> debug_info;


Any way to write this so that mutable isn't required. This can sometimes lead to race conditions.

yenong-amd · 2025-10-12T04:17:57Z

+
+  /// K dimension (reduction).
+  std::size_t k;
+


Can we have parameterized constructors for these types, e.g., dim3_t{size_t m_init, size_t n_init, size_t k_init)

Yes! We should do this as well.

NaveenElumalaiAMD · 2025-11-18T19:12:52Z

closing this PR for this one #2718

@neoblizz

## Technical Details Rebase @neoblizz (Muhammad) changes (#1859) with latest develop --------- Co-authored-by: Yen Ong <yen.ong@amd.com> Co-authored-by: Bryant Nelson <bryant.nelson@amd.com> Co-authored-by: neoblizz <osama94@gmail.com> Co-authored-by: yenong-amd <92903026+yenong-amd@users.noreply.github.com> Co-authored-by: Minsu Kim <minsu.kim@amd.com> Co-authored-by: Alex Brown <alex.brown@amd.com>

@ibrahimw1

## Testing Plan - Azure CI (deprecated): Runs all the Catch2 and new Python tests. ✅ - TheRock CI (integration): - CI workflow: https://github.com/ROCm/TheRock/actions/runs/20316557318 ✅ - TheRock branch: https://github.com/ROCm/TheRock/tree/users/neoblizz/minsukim-refactor ✅ (Some tests fail but build passed) - TheRock CI (libraries: hipblaslt, etc.) ✅ - Math CI (performance/hipblaslt) ✅ - Reviews from older PRs are addressed with the few exceptions. ✅ ## PRs History - Original PR: #1859 - Rebased PR: #2718 - Reverted PR: #3416 (comment) - Hot-fix PR: #3417 (review) ([TheRock CI Report](https://github.com/ROCm/TheRock/actions/runs/20293546191)) - **And now this!** ## Technical Details ### Refactor - [x] New file `types.hpp`, consolidates various origami types. - [x] New structs to replace the growing tuple. - [x] API updates to make it scalable, down from 20+ parameters to a few. - [x] Lots of redundant code removal. - [x] Reorganize into `types.hpp`; data types (see `hardware.hpp` for what needs to be moved) - [x] Refactor extract APIs - [x] Add enums for transpose - [x] Refactor debug/log reporting - [x] Remove mutable and statics - [x] Decouple `latency` out of `config_t` - [x] Rebase develop into this branch ### Python APIs & Unit Tests - [x] Update the tests - [x] Update rocroller - [x] hipblaslt/tensilelite/scripts to use the new API ### Testing Infrastructure - Replace YAML-based tests with Catch2 C++ tests - Replace GTest with Catch2 framework - Convert YAML-driven parameterized tests to pure C++ tests - Update common.hpp with direct C++ helper functions - Update CMakeLists.txt to use Catch2 instead of GTest - Remove Boost dependency (was only for YAML) - All test data now hardcoded in C++ for type safety - Better IDE support, debugging, and error messages ### Questions - Rocroller: coordinate on the API design, should they be reused? - Reuse dim3_t from hip_runtime.h (does HIP's dim3 have the same functionality)? - Should transpose/data-types be part of config as well for precomputing? ## Motivation This PR reapplies Origami project's refactor. 1. Make it simpler to make model updates. 2. Well-defined, scalable, clean interface. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. ## Work Moved to Future PRs - Logger APIs - Runtime Options - Move instruction map into separate file - Additional tests: #3301 - Minor batch-count specific changes: #3289 - Move Origami's Azure CI -> TheRock CI @ibrahimw1 --------- Co-authored-by: Brad Nemanich <Brad.Nemanich@amd.com> Co-authored-by: neoblizz <osama94@gmail.com>

Refactoring origami.

2c9617b

- New file , consolidates various origami types. - New structs to replace the growing tuple. - API updates to make it scalable, down from 20+ parameters to a few. - Lots of redundant code removal. '

github-actions Bot added project: hipblaslt shared: origami labels Sep 27, 2025

assistant-librarian Bot added the organization: ROCm label Sep 27, 2025

neoblizz added 3 commits September 26, 2025 20:40

Merge branch 'develop' into users/neoblizz/origami-refactor

09009e7

Formatting + Extract API.

e2ecfe9

Updating python to support new APIs.

90e0eb7

ryanswann-amd self-requested a review September 29, 2025 18:36

ryanswann-amd reviewed Sep 29, 2025

View reviewed changes

neoblizz added 3 commits September 29, 2025 15:41

Merge conflicts resolved w/ develop.

dcf06d5

Missing pragma once and header.

1ea1382

Refactor internal gemm implementation using new structs.

92ab844

neoblizz requested review from NaveenElumalaiAMD, pghysels, ryanswann-amd and yenong-amd September 30, 2025 04:45

Merge branch 'develop' into users/neoblizz/origami-refactor

298e271

neoblizz self-assigned this Sep 30, 2025

neoblizz requested review from AlexBrownAMD and awhittle3 September 30, 2025 15:46

bnemanich reviewed Sep 30, 2025

View reviewed changes

Comment thread shared/origami/include/origami/types.hpp Outdated

bnemanich reviewed Sep 30, 2025

View reviewed changes

neoblizz added 3 commits September 30, 2025 15:17

Removed hip-runtime dependency.

4a010dc

Addressing comments. Removes split from API.

6be7859

Merge branch 'develop' into users/neoblizz/origami-refactor

0078bd1

pghysels reviewed Oct 1, 2025

View reviewed changes

More clean-up.

0adbfba

bnemanich reviewed Oct 2, 2025

View reviewed changes

Switch to using Catch2 testing framework.

a06b069

neoblizz and others added 5 commits October 3, 2025 16:28

More comments addressed/cleanup.

5705f98

Remove old tests.

a5f1a88

Ensure header-only consumers get the same C++17 requirement.

46cc839

Support transpose_t enum.

04fe9f5

Decouple latency out of config_t

ac7c08e

yenong-amd reviewed Oct 12, 2025

View reviewed changes

Comment thread shared/origami/include/origami/hardware.hpp Outdated

ryanswann-amd and others added 7 commits October 15, 2025 14:42

Merge branch 'develop' into users/neoblizz/origami-refactor

0acff55

Merge branch 'develop' into users/neoblizz/origami-refactor

8c64418

Fix linker issue caused by std::execution::par call in origami

5ffe4b4

Update hipblaslt/tensilelite to use the refactored Origami APIs

68dbce1

Fix header file

ae7617d

Update rocroller to use the refactored origami APIs

4b1f6f3

Merge branch 'develop' into users/neoblizz/origami-refactor

f02dcc7

neoblizz assigned asunderwood Oct 22, 2025

Remove mutable and statics

3210fa0

NaveenElumalaiAMD mentioned this pull request Nov 17, 2025

[origami] Major refactoring of codebase #2718

Merged

15 tasks

NaveenElumalaiAMD closed this Nov 18, 2025

neoblizz mentioned this pull request Dec 17, 2025

[Origami] Reapply the origami refactor with fixes #3452

Merged

15 tasks

yenong-amd deleted the users/neoblizz/origami-refactor branch January 16, 2026 18:53

Conversation

neoblizz commented Sep 27, 2025 • edited by minsukim-amd Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Refactor

Python APIs & Unit Tests

Testing Infrastructure

Questions

Motivation

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

neoblizz Oct 1, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

NaveenElumalaiAMD commented Nov 18, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

8 participants

neoblizz commented Sep 27, 2025 •

edited by minsukim-amd

Loading

neoblizz Oct 1, 2025 •

edited

Loading