## Motivation
The grouped_gemm CK Tile kernel exists (e.g.,
`example/17_grouped_gemm/`) but has no Tile Engine wrapper. Grouped GEMM
handles multiple independent GEMM problems with varying M/N/K dimensions
in a single kernel launch. This PR adds the Tile Engine infrastructure
for automated kernel generation, benchmarking, and profiling of grouped
GEMM kernels.
Jira: AICK-809
## Technical Details
- Created Tile Engine wrapper under `tile_engine/ops/gemm/grouped_gemm/`
following the `gemm_universal` template
- Files added: `CMakeLists.txt`, `grouped_gemm_common.hpp`,
`grouped_gemm_benchmark.hpp`, `grouped_gemm_profiler.hpp`,
`grouped_gemm_benchmark.py`, `grouped_gemm_benchmark_single.cpp`,
`grouped_gemm_instance_builder.py`, `configs/`
- Supported datatypes: fp16, fp8, bf16, bf8
- Supported layouts: rcr, rrr, ccr, crr
- Target GPUs: gfx942, gfx950
- CK Tile kernel: `ck_tile::GroupedGemmKernel` from
`include/ck_tile/ops/gemm/kernel/grouped_gemm_kernel.hpp`
- Instance builder extends `GemmKernelBuilder` base class
- Registered in `tile_engine/ops/gemm/CMakeLists.txt`
- Updated Jenkinsfile to build and benchmark grouped_gemm targets in CI
- Benchmark infrastructure includes JSON output, CSV export, and
verification support
## Test Plan
- CMake configure succeeds for grouped_gemm targets
- Kernel instance builder generates valid kernel headers for all
(datatype, layout) combinations
- At least one kernel binary compiles and runs per datatype/layout
combination
- Correctness passes with `--verify 1` on gfx942/gfx950
## Test Result
## Submission Checklist
- [ ] 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: Copilot <175728472+Copilot@users.noreply.github.com>
## Motivation
Test grouped conv backward weight from ck tile and fix incorrect values.
## Technical Details
- Add test for CI
- Add daily tests
- Fix transpose load in BASE_V1 pipeline
## Test Plan
test_grouped_convnd_backward_weight_tile
## Test Result
in progress
## Submission Checklist
- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
AICK-783
## Motivation
While testing the Skip CI functionality, it revealed a minor issue where
the CI skip check fails when a branch is built at the exact commit where
it diverged from develop. CI is still run by default if a failure is
detected.
When git log <merge-base>..HEAD returns no files (because HEAD equals
merge-base), the command grep -v '^$' exits with error code 1, causing
the skip check to fail.
## Technical Details
Added || true to the grep commands so empty output is handled gracefully
instead of causing a script failure.
## Test Plan
- Simulate the failures and ensure the grep failure is handled
gracefully.
## Test Result
- Simulated grep failures using an empty string. The script handles the
error correctly.
- Verified the CI skip functionality skips CI when non-relevant file
changes are made.
- Verified the CI skip functionality does not skip CI when relevant file
changes are made.
## Submission Checklist
- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
## Motivation
Currently we have only fp8 and fp16 datatype support for universal GEMM
in Tile Engine with this PR support for bf8 and bf16 datatype will be
added during the CI phase
## Technical Details
Adding bf8 and bf16 support
## Test Plan
NA
## Test Result
NA
## Submission Checklist
- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
## Motivation
The CI skip logic has two issues that prevented it from working
correctly:
1. **Incorrect file patterns**: After migrating from standalone repo to
`rocm-libraries`, file paths now include the
`projects/composablekernel/` prefix (e.g.,
`projects/composablekernel/docs/README.md`). The skip patterns were
still checking for paths starting with `docs/`, which never matched.
2. **Incomplete build type support**: Jenkins multibranch pipelines
provide different environment variables for PR builds (`$CHANGE_TARGET`,
`$CHANGE_ID`) vs branch builds (`$BRANCH_NAME`). The previous logic only
compared `HEAD~1..HEAD` for branch builds, which missed changes from
multi-commit pushes and didn't properly handle feature branch builds.
When CI skipped or ran, there was no visibility into which files
triggered the decision, making it difficult to diagnose issues. You can
now see which files triggered the CI run.
## Technical Details
PR builds: Compares all commits against origin/$CHANGE_TARGET.
Feature branch builds: Uses git merge-base to find divergence point from
develop and checks all touched files since then.
Scheduled develop builds are unaffected. These builds are forced to run
from the pipeline parameters.
Example log output for PR Builds:
<img width="647" height="260" alt="image"
src="https://github.com/user-attachments/assets/c8673a81-acb2-4fb2-acbb-1c07b5ab3b69"
/>
Example log output for Branch Builds:
<img width="488" height="287" alt="image"
src="https://github.com/user-attachments/assets/fbb17ba7-eb2c-42a4-b820-b2a8b9e479c4"
/>
## Test Plan
Pre-PR validation (branch builds):
Push commits with only documentation changes → CI should skip. I will
have to verify this after this PR is merged!
Push commits with code changes → CI should run
Push commits that modify then revert code → CI should run (catching
reverts)
Verify debug output clearly shows skip/run decision
Post-PR validation (PR builds):
Create PR with only doc changes → CI should skip. I will have to verify
this after this PR is merged!
Create PR with mixed doc + code changes → CI should run and log which
files triggered it
Verify debug output clearly shows skip/run decision
## Test Result
All branch build checks succeeded.
All PR build checks succeeded.
## Submission Checklist
- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
## Motivation
Upgrade to official sccache version 0.14, since it now supports hip.
Also, switching daily builds from amd-staging to develop compiler
branch, since it should be more stable.
## Technical Details
<!-- Explain the changes along with any relevant GitHub links. -->
## Test Plan
<!-- Explain any relevant testing done to verify this PR. -->
## Test Result
<!-- Briefly summarize test outcomes. -->
## Submission Checklist
- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
## Motivation
This PR adds support for the gfx950 GPU architecture to the Tile Engine
in Composable Kernel library, focusing on GEMM operations with FP8 and
BF8 data types.
## Technical Details
Added gfx950-specific MFMA warp GEMM implementations with conditional
compilation.
Updated default GEMM configuration parameters for tile sizes and warp
configurations.
Added Jenkins CI pipeline stage for testing TILE_ENGINE_GEMM on gfx950
hardware.
## Test Plan
Tile engine itself is a benchmarking utility, so if it passes the CI it
will be tested automatically.
## Test Result
Tile engine itself is a benchmarking utility, so if it passes the CI it
will be tested automatically.
## Submission Checklist
- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
---------
Co-authored-by: Thrupti Raj Lakshmana Gowda<ThruptiRaj.LakshmanaGowda@amd.com>
Co-authored-by: Thomas Ning <Thomas.Ning@amd.com>
## Motivation
- ck mici jobs hanging at end, possibly at failure pattern checking
## Technical Details
- Disable failure pattern checking to see if hanging goes away
## Test Plan
- Observe behavior after merge
## Motivation
Fix the filter that determines whether CI builds are necessary.
## Technical Details
A script checks the files list returned by git diff and checks whether
any code source was modified. If not, if only documentation was changed,
it will allow skipping the builds. We make sure we only look at the
changes in projects/composablekernel/ folder.
## Motivation
Fixing some of the CK CI issues
## Technical Details
fixing paths to dockerfiles and scripts;
moving codegen tests to separate stage (collides with main build since
you must call cmake from same folder but different options);
fixing a couple of clang compilation issues with staging compiler;
## Motivation
- git diff needs access to reference repo
## Technical Details
- mount reference repo path into docker for selective_test_filter.py to
access
## Test Plan
- tested in MICI
## Test Result
- launch_tests.sh ran successfully
## Motivation
Pipelines were failing on Math CI status check.
## Technical Details
For the success case, we just changed the config in Jenkins to use a
proper app token and no code changes were required. However, the failure
case would not have worked as coded, so we needed to move that outside
of the `rocmnode()` block.
## Test Plan
I removed all of the CI in one of the commits to quickly test, and then
added it back. Got a successful "success" message and "failure" message
produced
## Motivation
- Corrects path to script due to superrepo migration
- Forces all tests to run by default
## Technical Details
- now in /projects/composablekernel
---------
Co-authored-by: illsilin_amdeng <Illia.Silin@amd.com>
## Motivation
- Maintain a reference repo on slave nodes that speeds up any
clone/checkout operations
## Technical Details
- clone a ref repo if it does not exist
- update ref repo if it does exist
- checkout after ref repo is updated
- eliminates double clone
## Test Result
- Initial checkouts succeeded
Implement per-page K/V quantization for paged attention:
- Add KV_BLOCKSCALE enum to BlockAttentionQuantScaleEnum
- Use exp2 shift trick to eliminate explicit P scaling overhead
- Prefetch physical pages offset for KV cache, overlaps with
computations
## Proposed changes
Please describe the motivation behind the pull request, whether it
enables a new feature or fixes a bug. If there are associated pull
requests or issues, please link them to the pull request.
## Checklist
Please put an `x` into the boxes that apply. You can also fill these out
after creating the PR. If you're not sure, please don't hesitate to ask.
- [ ] I have added tests relevant to the introduced functionality, and
the unit tests are passing locally
- [ ] I have added the test to REGRESSION_TESTS list defined at the top
of CMakeLists.txt in tests/CMakeLists.txt, **IF** the test takes more
than 30 seconds to run.
- [ ] I have added inline documentation which enables the maintainers
with understanding the motivation
- [ ] I have removed the stale documentation which is no longer relevant
after this pull request
- [ ] (If this change is user-facing) I have added release notes which
provide the end users with a brief summary of the improvement from this
pull request
- [ ] I have run `clang-format` on all changed files
- [ ] Any dependent changes have been merged
## Discussion
If this is a relatively large or complex change, feel free to start a
discussion by explaining why you chose the solution you did and what
alternatives you considered
---
🔁 Imported from
[ROCm/composable_kernel#3696](https://github.com/ROCm/composable_kernel/pull/3696)
🧑💻 Originally authored by @Jeff-Huang
---------
Co-authored-by: Jeff Huang <chiachi.huang@amd.com>
Co-authored-by: Illia Silin <Illia.Silin@amd.com>
## Motivation
Enable the CK CI after migration from standalone repo.
## Technical Details
Modify the jenkinsfile in projects/composablekernel to update the CI
workflow.
## Test Plan
This is for CK internal testing only.
## Test Result
Set up new CK CI pipeline/dashboard.
## Submission Checklist
- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
---------
Co-authored-by: Andrew Clark <andrew.clark@amd.com>
* Added two new failure patterns to detect. Including test function to verify if the patterns are detected
* Modifying pattern match to detect docker login failure. Removed passing tests.
* Removing passing tests. Modifying docker pattern to detect failure
* Removed passing tests
* Removing test logging function
[ROCm/composable_kernel commit: 421b714f13]
This change improves the clang-format CI check to be faster and not
depend on git being available in the build environment.
Changes:
- Use `find` instead of `git ls-files` (no git dependency)
- Check all C++ files: *.h, *.hpp, *.cpp, *.h.in, *.hpp.in, *.cpp.in, *.cl
- Exclude build/ and include/rapidjson directories
- Use parallel processing with 8 cores (-P 8) for ~8x speedup
- Show only errors with unified diff format (-u)
- Clear error messages: "ERROR: <file> needs formatting"
- Preserve original logic: run clang-format only when RUN_CPPCHECK=false,
or run both clang-format and cppcheck when RUN_CPPCHECK=true
Performance:
- Sequential processing: ~93 seconds for 5,899 files
- Parallel with 8 cores: ~12 seconds for 5,899 files
- Per-file processing time: ~15ms
This reduces CI time while maintaining code formatting standards.
[ROCm/composable_kernel commit: 98abfa4ade]
1. Added `-DCK_EXPERIMENTAL_BUILDER=OFF` to the `setup_args` to explicitly disable the experimental builder
2. Added a detailed comment explaining why this is necessary:
- SLES15 is a legacy platform with limited C++20 ecosystem support
- While the ROCm compiler supports C++20, the older system libraries and standard library implementation on SLES15 does not reliably support all C++20 features required by the experimental CK Builder
[ROCm/composable_kernel commit: 2d233c838a]
* memory op changes
* memory op changes
* Fixing TILE_ENGINE_BASIC in Tile Engine
* Removing gfx90a from Tile Engine Run
* [CK TILE ENGINE] increasing ci configs for BASIC case
* Setting RUN_TILE_ENGINE_BASIC_TESTS to ON by default
---------
Co-authored-by: Max Podkorytov <4273004+tenpercent@users.noreply.github.com>
[ROCm/composable_kernel commit: 51027474af]
* solve compiler issue
* solve the gfx950 mfma shuffle regression
* refactor jenkinsfile to handle arch name better
* [CK TILE] set divisor to count of thread along k dimension
* fix the compiler error
* solve degradation
* Finish the multiplies fix
* fix the scales
* solve compilation error
* solve the composes
* solve the error of tile sweeper
* fix the test and example
* fix for gfx950
---------
Co-authored-by: Max Podkorytov <4273004+tenpercent@users.noreply.github.com>
Co-authored-by: illsilin_amdeng <Illia.Silin@amd.com>
Co-authored-by: Cong Ma <congma13@amd.com>
[ROCm/composable_kernel commit: 00c46785a8]
* Removing hard-coded trace filename
* Including stage name in notification
* Simplifying capture setup and tagging file names with arch
* Removed test property from notification message
* Fixing regex to get arch name
* Fixing error in notification and modified regex
[ROCm/composable_kernel commit: e77a7ca2bc]
* generate and visualize build traces for all archs
* generate build traces in all cases
* fix jenkins logic
* fix typo
* use more threads for parsing dependency map
* add script to parse ninja traces and issue warnings
* fix python script syntax and header
* fix python syntax one more time
* fix python syntax
[ROCm/composable_kernel commit: 3dfa794fab]
* Parallelization in dataset generation
* Parallelizable tests for fwd, bwd data, bwd weight with datasets
* .gitignore generated datasets
* Test parallelization script with round-robin GPU scheduling
* Parallelization updates to test generation and running
* Dataset paths relative to executable
* Update output from test generation
* Default to one GPU in test generation
* Add small dataset tests to Jenkins
* Update copyright lines
* Update test_data/generate_test_dataset.sh
Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
* Move trap disable
* Common get path function
---------
Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
[ROCm/composable_kernel commit: fe35ba5dac]
