Backport CK develop onto release/rocm-rel-7.0.2.1 for the ROCm 7.0.2.1 delivery

.ck-project-root

@@ -0,0 +1,2 @@
+# Composable Kernel project root marker
+# Used by scripts in script/tools/ to locate the project root

.github/CODEOWNERS

@@ -1,8 +1,8 @@
-* @illsilin @carlushuang @qianfengz @aosewski @poyenc @geyyer @bartekxk @andriy-ca @afagaj @asleepzzz @tenpercent @ThomasNing @coderfeli @shumway @vidyasagar-amd
+* @illsilin @carlushuang @qianfengz @aosewski @poyenc @geyyer @bartekxk @andriy-ca @afagaj @asleepzzz @tenpercent @coderfeli @cgmillette @shumway @vidyasagar-amd @vpietila-amd @Snektron @JiaLuo-CAN
 # Documentation files
-docs/ @ROCm/rocm-documentation @illsilin @carlushuang @qianfengz @aosewski @poyenc @geyyer @bartekxk @andriy-ca @afagaj @asleepzzz @ThomasNing @coderfeli @shumway @vidyasagar-amd @ddembeckAMD
-*.md @ROCm/rocm-documentation @illsilin @carlushuang @qianfengz @aosewski @poyenc @geyyer @bartekxk @andriy-ca @afagaj @asleepzzz @ThomasNing @coderfeli @shumway @vidyasagar-amd @ddembeckAMD
-*.rst @ROCm/rocm-documentation @illsilin @carlushuang @qianfengz @aosewski @poyenc @geyyer @bartekxk @andriy-ca @afagaj @asleepzzz @ThomasNing @coderfeli @shumway @vidyasagar-amd @ddembeckAMD
-.readthedocs.yaml @ROCm/rocm-documentation @illsilin @carlushuang @qianfengz @aosewski @poyenc @geyyer @bartekxk @andriy-ca @afagaj @asleepzzz @ThomasNing @coderfeli @shumway @vidyasagar-amd @ddembeckAMD
+docs/ @ROCm/rocm-documentation @illsilin @carlushuang @qianfengz @aosewski @poyenc @geyyer @bartekxk @andriy-ca @afagaj @asleepzzz @coderfeli @cgmillette @shumway @vidyasagar-amd @ddembeckAMD @vpietila-amd @Snektron @JiaLuo-CAN
+*.md @ROCm/rocm-documentation @illsilin @carlushuang @qianfengz @aosewski @poyenc @geyyer @bartekxk @andriy-ca @afagaj @asleepzzz @coderfeli @cgmillette @shumway @vidyasagar-amd @ddembeckAMD @vpietila-amd @Snektron @JiaLuo-CAN
+*.rst @ROCm/rocm-documentation @illsilin @carlushuang @qianfengz @aosewski @poyenc @geyyer @bartekxk @andriy-ca @afagaj @asleepzzz @coderfeli @cgmillette @shumway @vidyasagar-amd @ddembeckAMD @vpietila-amd @Snektron @JiaLuo-CAN
+.readthedocs.yaml @ROCm/rocm-documentation @illsilin @carlushuang @qianfengz @aosewski @poyenc @geyyer @bartekxk @andriy-ca @afagaj @asleepzzz @coderfeli @cgmillette @shumway @vidyasagar-amd @ddembeckAMD @vpietila-amd @Snektron @JiaLuo-CAN
 # Header directory for Doxygen documentation
-library/include/ @ROCm/rocm-documentation @illsilin @carlushuang @qianfengz @aosewski @poyenc @geyyer @bartekxk @andriy-ca @afagaj @asleepzzz @ThomasNing @coderfeli @shumway @vidyasagar-amd
+library/include/ @ROCm/rocm-documentation @illsilin @carlushuang @qianfengz @aosewski @poyenc @geyyer @bartekxk @andriy-ca @afagaj @asleepzzz @coderfeli @cgmillette @shumway @vidyasagar-amd @vpietila-amd @Snektron @JiaLuo-CAN

.gitignore

@@ -36,7 +36,13 @@ tags
 # Editors
 .vscode
 
-# build-in-source directory
+# CMake formatting configuration (local)
+.cmake-format.yaml
+
+# Cline
+.cline*
+
+# build-in-source directory (see exceptions below)
 build*
 
 # emacs temporary/backup files
@@ -58,16 +64,75 @@ _doxygen/
 docs/doxygen/html
 docs/doxygen/xml
 
-# JetBrains IDE
+# JetBrains IDE (see build* exceptions below)
 .idea/
 cmake-build*/
 build*/
 
+# LSP configuration
+.clangd
+
+# User-defined CMake presets
+CMakeUserPresets.json
+
 # Python virtualenv
 .venv/
 
 # Python cache
 __pycache__/
 
+# Cache directories
 .cache/
-
+.ck_tile_cache/
+ck_tile_cache/
+**/kernel_cache/
+**/.kernel_cache/
+
+# Dispatcher kernel cache (user-generated, can be large)
+dispatcher/**/kernel_cache/
+dispatcher/**/.kernel_cache/
+dispatcher/**/cached_kernels/
+dispatcher/**/*.hsaco
+dispatcher/**/*.co
+
+# Dispatcher generated JSON exports
+dispatcher/**/*_kernels.json
+dispatcher/**/dispatcher_kernels.json
+
+# Generated test data
+test_data/*
+!test_data/*.py
+!test_data/*.sh
+!test_data/requirements.txt
+
+# Exceptions to build* patterns above
+# The experimental/builder directory should be tracked despite matching build*
+!experimental/builder
+!experimental/builder/**
+# Heuristics: benchmark data (never in git)
+dispatcher/heuristics/data/
+
+# Heuristics: experimental/training artifacts (exclude from git)
+dispatcher/heuristics/models/**/oof_predictions.parquet
+dispatcher/heuristics/models/**/cv_metrics_*.json
+dispatcher/heuristics/models/**/eval_report.json
+dispatcher/heuristics/models/**/feature_importances_*.json
+dispatcher/heuristics/models/**/model_tflops_ihem.lgbm
+dispatcher/heuristics/models/**/model_tflops_log.lgbm
+dispatcher/heuristics/models/**/model_tflops_log_big.lgbm
+
+# Heuristics: keep in git (production model files):
+#   models/{op}_{dtype}_{arch}/model_tflops.lgbm
+#   models/{op}_{dtype}_{arch}/model_latency.lgbm
+#   models/{op}_{dtype}_{arch}/model_bandwidth.lgbm
+#   models/{op}_{dtype}_{arch}/feature_spec.json
+#   models/{op}_{dtype}_{arch}/train_manifest.json
+
+# Heuristics: logs and caches
+dispatcher/heuristics/*.log
+dispatcher/heuristics/__pycache__/
+dispatcher/heuristics/tests/__pycache__/
+dispatcher/heuristics/.pytest_cache/
+
+# CLAUDE markdown files
+CLAUDE.md

.pre-commit-config.yaml

@@ -1,38 +1,53 @@
 repos:
--   repo: local
+-   repo: https://github.com/pre-commit/mirrors-clang-format
+    rev: v18.1.3
     hooks:
     -   id: clang-format
-        name: clang-format
-        entry: clang-format-12 -i --style=file
-        language: system
         types_or: [c++, inc]
-    -   id: copyright-year-checker
-        name: copyright-year-checker
-        entry: script/check_copyright_year.sh
+-   repo: https://github.com/astral-sh/ruff-pre-commit
+    rev: v0.14.0
+    hooks:
+    -   id: ruff-check
+        args: [ --fix ]
+        exclude: |
+            (?x)^(
+                docs/conf.py
+            )$
+    -   id: ruff-format
+        exclude: |
+            (?x)^(
+                docs/conf.py
+            )$
+-   repo: local
+    hooks:
+    -   id: copyright-header-checker
+        name: Check copyright headers
+        entry: projects/composablekernel/script/check_copyright_year.sh
         verbose: false
         language: script
-        types: [c++]
+        types_or: [c++, python, shell, cmake]
+    -   id: ascii-only-checker
+        name: Check for non-ASCII characters in C/C++ sources
+        entry: projects/composablekernel/script/check_ascii_only.sh
+        language: script
+        types_or: [c++, inc]
+    -   id: crlf-checker
+        name: Check for CRLF line endings in C/C++ sources
+        entry: projects/composablekernel/script/check_no_crlf.sh
+        language: script
+        types_or: [c++, inc]
     -   id: remove-exec-bit
         name: Remove executable bit from non-executable files
-        entry: script/remove_exec_bit.sh
+        entry: projects/composablekernel/script/remove_exec_bit.sh
         language: script
         types_or: [c++, text]
         verbose: true
-    -   id: ruff-check
-        name: Ruff Linter
-        entry: ruff check --fix
+    -   id: remod-ck-tile
+        name: Run ck_tile remod.py
+        entry: python projects/composablekernel/script/remod_for_ck_tile.py
         language: python
-        types: [python]
-        additional_dependencies: [ruff]
-    -   id: ruff-format
-        name: Ruff Formatter
-        entry: ruff format
-        language: python
-        types: [python]
-        additional_dependencies: [ruff]
-    -   id: run-remod-if-ck-tile-changed
-        name: Run remod.py if ck_tile files changed
-        entry: script/remod_for_ck_tile.sh
-        language: script
-        always_run: true
+        files: '^projects/composablekernel/(include|example)/ck_tile/.*$'
+        additional_dependencies:
+        - dos2unix
+        - clang-format==18.1.3
         pass_filenames: false

.readthedocs.yaml

@@ -4,13 +4,13 @@
 version: 2
 
 sphinx:
-   configuration: docs/conf.py
+   configuration: projects/composablekernel/docs/conf.py
 
 formats: [htmlzip, pdf, epub]
 
 python:
    install:
-   - requirements: docs/sphinx/requirements.txt
+   - requirements: projects/composablekernel/docs/sphinx/requirements.txt
 
 build:
    os: ubuntu-22.04

ACRONYMS.md

@@ -0,0 +1,67 @@
+# Acronyms in Composable Kernel
+
+The following acronyms are used in the Composable Kernel codebase:
+
+| Acronym | Expansion | Explanation |
+|---------|-----------|-------------|
+| BF16    | Brain Floating Point 16 | 1 Signed bit, 8 Exponent bits, 7 Significand bits |
+| BF8     | 8-bit Brain Floating Point | 1 Signed bit, 3 Exponent bits, 4 Significand bits |
+| DLA     | Deep Learning Accelerator | Specialized hardware for deep learning workloads |
+| DRAM    | Dynamic Random-Access Memory | Main memory. Global memory on GPU |
+| E2E     | End-to-End | Complete pipeline or process from input to output |
+| ELU     | Exponential Linear Unit | Activation function: $x$ if $x>0$ else $\alpha(e^x-1)$ |
+| FMHA    | Fused Multi-Head Attention | Efficient transformer attention kernel, fusing softmax, masking, and matmul |
+| FP16    | Half-Precision Floating Point | 16-bit IEEE floating point format |
+| FP32    | Single-Precision Floating Point | 32-bit IEEE floating point format |
+| FP64    | Double-Precision Floating Point | 64-bit IEEE floating point format |
+| FP8     | 8-bit Floating Point | Experimental 8-bit floating point format for inference |
+| GEMM    | General Matrix Multiply | Matrix multiplication operation: $C = A \times B$ |
+| GELU    | Gaussian Error Linear Unit | Activation function: $x \cdot \Phi(x)$ |
+| GQA     | Grouped Query Attention | Variant of multi-head attention with grouped queries/keys/values |
+| HBM     | High Bandwidth Memory | Fast memory used in modern GPUs |
+| HIP     | Heterogeneous-Compute Interface for Portability | AMD's CUDA-like GPU programming API |
+| INT8    | 8-bit Integer | Quantized integer format for inference |
+| KVS     | Key-Value Store | Data structure for storing key-value pairs (context: QKV in transformers) |
+| L2/L1   | Level 2/Level 1 Cache | On-chip memory hierarchy in CPUs/GPUs |
+| LDS     | Local Data Share | Shared memory on AMD GPUs (equivalent to CUDA's shared memory) |
+| LLM     | Large Language Model | Transformer-based model for NLP tasks |
+| LSE     | Log-Sum-Exp | Numerically stable softmax computation: $\log(\sum \exp(x))$ |
+| MHA     | Multi-Head Attention | Attention mechanism with multiple heads in transformers |
+| MFMA    | Matrix Fused Multiply-Add | AMD GPU hardware instruction for matrix-matrix multiplication |
+| MoE     | Mixture of Experts | Neural network architecture with multiple expert subnetworks |
+| MQA     | Multi-Query Attention | Variant of multi-head attention with shared keys/values across heads |
+| RCCL    | ROCm Collective Communications Library | AMD Library for multi-GPU communication |
+| NCHW    | Batch, Channel, Height, Width | Tensor layout: batch-major, channels-first |
+| NHWC    | Batch, Height, Width, Channel | Tensor layout: batch-major, channels-last |
+| OOM     | Out Of Memory | Error when memory allocation fails |
+| QAT     | Quantization Aware Training | Training technique for quantized inference |
+| QKV     | Query, Key, Value | Components of transformer attention mechanism |
+| RDMA    | Remote Direct Memory Access | High-speed network memory access |
+| RDQuant | Rowwise Dynamic Quantization | Quantization technique with per-row scaling for int8 inference |
+| ReLU    | Rectified Linear Unit | Activation function: $\max(0, x)$ |
+| ROCm    | Radeon Open Compute | AMD's open GPU computing stack |
+| SGD     | Stochastic Gradient Descent | Optimization algorithm for training neural networks |
+| SM      | Streaming Multiprocessor | GPU compute unit (NVIDIA terminology) |
+| SWA     | Sliding Window Attention | Attention mechanism with a limited window for each token |
+| TLB     | Translation Lookaside Buffer | Memory management unit cache for virtual-to-physical address translation |
+| VGPR    | Vector General Purpose Register | GPU register for vector operations |
+| WARP    | Group of Threads | Smallest scheduling unit on NVIDIA GPUs (32 threads) |
+| WMMA    | Warp Matrix Multiply-Accumulate | NVIDIA's matrix-multiply hardware primitive |
+| XLA     | Accelerated Linear Algebra | Compiler for optimizing ML computations (Google) |
+
+### Common Variable Acronyms in Code
+
+| Symbol | Meaning | Context |
+|--------|---------|---------|
+| M, N, K | Matrix dimensions | GEMM: $A[M,K] \times B[K,N] = C[M,N]$ |
+| Q, K, V | Query, Key, Value | Transformer attention |
+| S       | Sequence length | NLP, transformers |
+| D       | Dimension | Hidden size, feature dim |
+| B       | Batch size | ML batch processing |
+| H       | Head count | Multi-head attention |
+| C       | Channel | CNNs, tensor layouts |
+| T       | Token | NLP, sequence models |
+
+---
+
+If you find an acronym not listed here, please submit a pull request or issue!

CHANGELOG.md

@@ -2,6 +2,88 @@
 
 Documentation for Composable Kernel available at [https://rocm.docs.amd.com/projects/composable_kernel/en/latest/](https://rocm.docs.amd.com/projects/composable_kernel/en/latest/).
 
+## Composable Kernel 1.2.0 for ROCm 7.13
+
+### Added
+
+* Added overload of load_tile_transpose that takes reference to output tensor as output parameter
+* Use data type from LDS tensor view when determining tile distribution for transpose in the GEMM pipeline
+* Added eightwarps support for abquant mode in blockscale GEMM.
+* Added preshuffleB support for abquant mode in blockscale GEMM.
+* Added support for explicit GEMM in CK_TILE grouped convolution forward and backward weight.
+* Added TF32 convolution support on gfx942 and gfx950 in CK. It could be enabled/disabled via `DTYPES` of "tf32".
+* Added streamingllm sink support for FMHA FWD, include qr_ks_vs, qr_async and splitkv pipelines.
+* Added support for microscaling (MX) FP8/FP4 mixed data types to Flatmm pipeline.
+* Added support for fp8 dynamic tensor-wise quantization of fp8 fmha fwd kernel.
+* Added FP8 KV cache support for FMHA batch prefill.
+* Added support for gfx1153 target.
+* Added FMHA batch prefill kernel support for several KV cache layouts, flexible page sizes, and different lookup table configurations.
+* Added gpt-oss sink support for FMHA FWD, include qr_ks_vs, qr_async, qr_async_trload and splitkv pipelines.
+* Added persistent async input scheduler for CK Tile universal GEMM kernels to support asynchronous input streaming.
+* Added FP8 block scale quantization for FMHA forward kernel.
+* Added gfx11 support for FMHA.
+* Added microscaling (MX) FP8/FP4 support on gfx950 for FMHA forward kernel ("qr" pipeline only).
+* Added FP8 per-tensor quantization support for FMHA forward V3 pipeline on gfx950.
+
+### Changed
+
+### Upcoming changes
+
+## Composable Kernel 1.2.0 for ROCm 7.2.0
+
+### Added
+* Added tests for f8 x bf8 on CompV3, and f8 x bf8 with K_BlockSize 32 on CompV4
+* Added CK-Tile dispatcher - a unified kernel dispatch, code generation and architecture-based kernel filtering system with with C++ and Python frontends starting with GEMM support.
+* Added support for bf16 data type to grouped_gemm and grouped_gemm_preshuffle.
+* Added Col-Col-Row-Col layout support for aquant mode in blockscale GEMM.
+* Added support for mixed precision fp8 x bf8 universal GEMM and weight preshuffle GEMM.
+* Added a compute async pipeline in the CK Tile universal GEMM on gfx950.
+* Added support for B Tensor type `pk_int4_t` in the CK Tile weight preshuffle GEMM.
+* Added the new api to load different memory sizes to SGPR.
+* Added support for B Tensor preshuffle in CK Tile grouped GEMM.
+* Added a basic copy kernel example and supporting documentation for new CK Tile developers.
+* Added support for grouped GEMM kernels to perform Multi D elementwise operation.
+* Added support for multiple ABD GEMM.
+* Added benchmarking support for tile engine GEMM Multi D.
+* Added block scaling support in CK Tile GEMM, allowing flexible use of quantization matrices from either A or B operands.
+* Added the row-wise column-wise quantization for CK Tile GEMM and CK Tile grouped GEMM.
+* Added support for f32 to FMHA (forward and backward).
+* Added tensor-wise quantization for CK Tile GEMM.
+* Added support for batched contraction kernel.
+* Added WMMA (gfx12) support for FMHA.
+* Added pooling kernel in CK_TILE
+* Added top-k sigmoid kernel in CK_TILE
+* Added the blockscale 2D support for CK_TILE GEMM.
+* Added Flatmm pipeline for microscaling (MX) FP8/FP4 data types
+* Added reduce and multi reduction kernels
+
+### Changed
+
+* Removed `BlockSize` in `make_kernel` and `CShuffleEpilogueProblem` to support Wave32 in CK Tile (#2594)
+* Added an optional template parameter `Arch` (`gfx9_t`, `gfx12_t` etc.) to `make_kernel` to support linking multiple object files that have the same kernel compiled for different architectures.
+* FMHA examples and tests can be built for multiple architectures (gfx9, gfx950, gfx12) at the same time.
+
+### Upcoming changes
+
+* Composable Kernel will be adopting C++20 features in an upcoming ROCm release, updating the minimum compiler requirement to C++20. Ensure that your development environment complies with this requirement to facilitate a seamless transition.
+
+## Composable Kernel 1.1.0 for ROCm 7.1.1
+
+### Upcoming changes
+
+* Composable Kernel will be adopting C++20 features in an upcoming ROCm release, updating the minimum compiler requirement to C++20. Ensure that your development environment complies with this requirement to facilitate a seamless transition.
+
+## Composable Kernel 1.1.0 for ROCm 7.1.0
+
+### Added
+
+* Added support for hdim as a multiple of 32 for FMHA (fwd/fwd_splitkv/bwd)
+* Added support for elementwise kernel.
+
+### Upcoming changes
+
+* Non-grouped convolutions are deprecated. Their functionality is supported by grouped convolution.
+
 ## Composable Kernel 1.1.0 for ROCm 7.0.0
 
 ### Added
@@ -23,20 +105,15 @@ Documentation for Composable Kernel available at [https://rocm.docs.amd.com/proj
 * Added Ping-pong scheduler support for GEMM operation along the K dimension.
 * Added rotating buffer feature for CK_Tile GEMM.
 * Added int8 support for CK_TILE GEMM.
+* Added CK Tile Epilogue Chainer framework for composable epilogue sequences in GEMM operations
 
 ### Optimized
 
+* Optimize the gemm multiply multiply preshuffle & lds bypass with Pack of KGroup and better instruction layout.
+* Added Vectorize Transpose optimization for CK Tile
+* Added the asynchronous copy for gfx950
 
-* Optimize the gemm multiply multiply preshuffle & lds bypass with Pack of KGroup and better instruction layout. (#2166)
-* Added Vectorize Transpose optimization for CK Tile (#2131)
-* Added the asynchronous copy for gfx950 (#2425)
-
-
-### Fixes
-
-None
-
-### Changes
+### Changed
 
 * Removed support for gfx940 and gfx941 targets (#1944)
 * Replaced the raw buffer load/store intrinsics with Clang20 built-ins (#1876)
@@ -45,10 +122,6 @@ None
 * Number of instances in instance factory for grouped convolution backward weight NGCHW/GKYXC/NGKHW has been reduced.
 * Number of instances in instance factory for grouped convolution backward data NGCHW/GKYXC/NGKHW has been reduced.
 
-### Known issues
-
-None
-
 ## Composable Kernel 1.1.0 for ROCm 6.1.0
 
 ### Additions