7 Commits

Author SHA1 Message Date
Christopher Millette
d7836ff0b2 [rocm-libraries] ROCm/rocm-libraries#5222 (commit 4fe0911)
[CK_TILE] Fix MMA layout test to match amdgcn_mma OpFamily
 parameter (#5222)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

## Summary

- PR #4837 added `MmaOpFamily OpFamily_` as a new template parameter to
`amdgcn_mma` and `MmaDefaultSelector`, but the MMA layout test (PR
#4495) was not updated to include it
- Add the missing `OpFamily_` parameter to all three `RegisterMapTraits`
partial specializations (gfx9, gfx11, gfx12) and all
`MmaDefaultSelector` usages
- Fixes build failure: `template argument for non-type template
parameter must be an expression`

## Test plan

- [x] Verified test compiles cleanly with ROCm 7.1.1 clang++ targeting
gfx90a
- [x] `test_amdgcn_mma_layout` gfx90a (MFMA): PASSED
- [x] `test_amdgcn_mma_layout` gfx1201 (WMMA): SKIPPED (no device)
- [x] `test_amdgcn_mma_layout` gfx1100 (WMMA): SKIPPED (no device)
- [x] CI validation on all GPU targets

🤖 Generated with [Claude Code](https://claude.com/claude-code)
2026-03-09 05:18:50 +00:00
Wojciech Laskowski
b80e41f3bc [rocm-libraries] ROCm/rocm-libraries#4495 (commit 5664eb0)
Adding layout test for amdgcn_mma structs

## Motivation

Currently, the test suite for `amdgcn_mma` focuses on the design (e.g.
choosing the correct specialization based on SFINAE) and a single live
test that checks if selected MmaOp runs. This PR adds a simplified GEMM
test kernel that checks the exact layout of the selected MmaOp.

## Technical Details

The test in `test_amdgcn_mma_layout.cpp` launches MxKxN test cases (one
per block), where each case:
1. Constructs A and B tensors on a device with a single 1 at A(m,k) and
B(k,n) (rest is all 0s)
  2. Executes the MMA intrinsic.
  3.  Checks if C has the "1" on the excpeted position.

For the MMA instrinsic, it pulls a Mma op from amdgcn_mma specialization
based on a given input (tile dimension, data types).

Note 1: As a helper, in `test_amdgcn_mma_layout_util.hpp` we add
register map for a given amdgcn_mma specialization. Register mapping is
currently based on the `tile_distribution_encoding`.

Note 2: Everything is added to the test suite, no additions to the
actual `amdgcn_mma` structs. All the extra information that is needed,
but not yet provided by `amdgcn_mma` structs, is added as a boilerplate
to the header. TODO: Rebase this PR on top of the `amdgcn_mma` refactor
or clean it up after merge.

## Test Plan

This PR solely adds a new test to the existing code.

## Test Result

Tests pass.

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
2026-03-06 16:22:16 +00:00
chris-tsiaousis-hpc
03ce21ddcb [rocm-libraries] ROCm/rocm-libraries#4837 (commit 6316035)
[CK TILE] Unification of sparse MFMA/WMMA policy structs
 (#4837)

## Motivation

The existing unification work supports DENSE intrinsics. In this PR we
enable support for SPARSE as well as SCALE intrinsics and add an example
SPARSE implementation.

## Technical Details

Mostly trivial changes. One framework change is that the desired
`MmaOpFamily` is passed to the `MmaDefaultSelector`. As my relevant
commit explains, we do not support a fallback family at the moment, but
it is something we can consider.

## Test Plan

Added a new test for the relevant sparse specializations.

## Test Result

Test should pass.

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
2026-03-05 19:53:16 +00:00
Cong Ma
3af1a0aafc [rocm-libraries] ROCm/rocm-libraries#4355 (commit e7f6909)
[CK TILE] Refactor sequence_reverse_inclusive_scan

## Proposed changes

Refactor ck tile `sequence_reverse_inclusive_scan` from recursive to
for-loop.

Tracking issue: #4229

This pull request introduces a new lightweight array type,
`static_array`, and refactors the sequence utilities to use it for
improved constexpr support and simplicity. The changes also include
updates to the build system to add container-related tests.

**Core Library Improvements:**

* Added a new header `static_array.hpp` that defines the `static_array`
type, a constexpr-friendly array with basic accessors and no custom
constructors.
* Updated includes in `core.hpp` and `sequence.hpp` to import
`static_array`.
[[1]](diffhunk://#diff-14b406eccf59794051a16c0c9c1a7e11234324bfdd107a5bbe0f173cd25bcddcR44)
[[2]](diffhunk://#diff-5042e5b47bb2ba78bbab2d284338cf0503bc8fb76a7d631cc2684ad6ca832a76R7)

**Refactoring to Use `static_array`:**

* Refactored sequence utilities in `sequence.hpp` to use `static_array`
instead of the previously forward-declared `array` type, including in
histogram and array generation logic.
[[1]](diffhunk://#diff-5042e5b47bb2ba78bbab2d284338cf0503bc8fb76a7d631cc2684ad6ca832a76L1108-R1133)
[[2]](diffhunk://#diff-5042e5b47bb2ba78bbab2d284338cf0503bc8fb76a7d631cc2684ad6ca832a76L1130-R1146)
* Rewrote the implementation of `sequence_reverse_inclusive_scan` to use
`static_array` for intermediate storage, improving constexpr evaluation
and clarity.

**Build System and Testing:**

* Added a new test subdirectory for container tests and a GoogleTest
executable for `unit_sequence.cpp` to the CMake build configuration.
[[1]](diffhunk://#diff-5d35ff7555d3f0b438d45cde06b661eb1332cdbec66287ac7ec3c478d688aae5R5)
[[2]](diffhunk://#diff-1f54f0d2b431b7fc74f7b4ffb66e80c381c904c3383b1d27987467e3482d6d7aR1-R7)

Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
2026-02-24 15:52:33 +00:00
Aviral Goel
6d25525adc feat(precommit-hooks): add check for correct copyright header (#3302)
* chore(copyright): update copyright header for left files

* feat(copyright): add copyright check to precommit hooks

* chore(copyright): update copyright header for include/ck_tile directory

* chore(copyright): update copyright header for example directory

* chore(copyright): update copyright header for .github directory

* refactor: copyright_check script with better if else handling

* chore(copyright): update compyright header for remaining files

* feat: add script to automate copyright addition
2025-12-10 22:50:43 -08:00
Aviral Goel
004784ef98 chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
* chore(copyright) update library wide CMakeLists.txt files copyright header template

* Fix build

---------

Co-authored-by: Sami Remes <samremes@amd.com>
2025-11-28 13:49:54 -08:00
Christopher Millette
b9c6cb1452 First look at mfma / wmma unification (#2704)
* First look at mfma / wmma unification

* Refactor

* Re-org file structure

* Restructure transform selection and WaveWiseMma class

* Update license files. Add missing gfx1151 support. Change wave size for HOST to 1. Update datatypes naming consistency

* Fixes default MmaSelector implentation

* Adds unit tests for amdgcn_mma and arch

* Consolidate common arch id checks to constexpr functions. Strongly type ids as amdgcn_target_arch_id object.

* Refactor is_any_value_of

* Fixes mma_selector logic

* Fix typo

* Add mma selector test for tile decomposition

* Fix compilation of mma.hpp

* Revert back to c++17 compatibility

* Fix compiler error by returning index_t from get_warp_size()

* Apply suggestions from code review

Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>

* Fixes compiler error for missing is_wave32() function

* Fixes compiler error for host wave_size() should be 64

* Fixes compiler errors where __cpp_concepts is not defined

* Fixes compiler errors where __cpp_concepts is not defined

* Fix test failure for host is wave64 by default

---------

Co-authored-by: Chris Millette <you@example.com>
Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
2025-11-24 09:39:59 -08:00