Commit Graph

16 Commits

Author SHA1 Message Date
Casey-Shi
3bcbdd608e [Tile Engine] Add benchmark for tile engine gemm. (#2193)
* initial commit -m benchmark

* only support profile

* fix

* fix doc

* add default config

* add ci

* fix cmake

* tmp save for gen blobs

* fix bug

* merge

* range config

* test success

* fix

* fix

* move struct

* remove config property

* fix config

* remove comment

* add cmake option & modify

* add changelog

* fix

* format

* add pydantic module to the docker image

* fix

* add benchmark for cold and warmp up

* python format

* add asm cache control

* fix README

* remove pydantic module

* modify changelog

* fix config

* recover benchmark_gemm and fix

* format python

* refactor profiler

* fix csv bug

* fix codegen bug

* add kernel instance object

* add benchmark gemm executable

* fix jenkins & delete extra header

* disable warning output & enable default config

* Disable sparsity for invalid warp tile combinations

* fix gemm host template func

* refactor gemm profiler

* filter out some inmstances

* default config test & fix codegen bug

* add sparse flag to gen more instances

---------

Co-authored-by: illsilin <Illia.Silin@amd.com>
Co-authored-by: khuagarw <khuagarw@amd.com>
Co-authored-by: Thomas Ning <Thomas.Ning@amd.com>

[ROCm/composable_kernel commit: 128f5a1eab]
2025-05-26 22:32:36 -07:00
Khushbu Agarwal
395262f196 Adding validation for tile sizes in Tile Engine (#2189)
* Adding validation for tile sizes

* Add architecture in config, and shuffle lines of code in warp_gemm.hpp

* Enable MFMA for gfx950, and invalid tile handling

[ROCm/composable_kernel commit: 3d8d6e75e4]
2025-05-15 10:28:31 -07:00
Khushbu Agarwal
699091846a Disable SMFMA gfx90a (#2184)
* sparsity fix for gfx90a

* reverting tile_engine changes

[ROCm/composable_kernel commit: f05e45ba59]
2025-05-12 09:56:23 -07:00
Khushbu Agarwal
83180829b6 Disable SMFMA for gfx90a (#2182)
[ROCm/composable_kernel commit: ef72a4b9bc]
2025-05-09 00:18:07 -07:00
Thomas Ning
9fa7eed532 Revert "Disable the SMFMA instruction for gfx90a. (#2174)" (#2175)
This reverts commit 95b372278a9a57ac2be99700105674f25e5a3554.

[ROCm/composable_kernel commit: c757046d49]
2025-05-08 00:07:03 -07:00
Khushbu Agarwal
f7db78cfcf Disable the SMFMA instruction for gfx90a. (#2174)
* remove smfma for gfx90a

* clang formatted

[ROCm/composable_kernel commit: a32d907771]
2025-05-07 23:09:22 -07:00
Khushbu Agarwal
a86620111c [CK_Tile] Simplified Mem pipeline (#2159)
* simplify code

* compiled the code

* Simplified example and codegen for mem pipeline

* Reveting config and universal gemm example

* clang formatted

* remove comments

* clang formatted

* Add memory operation changes for defualt pipeline

* fix config file

---------

Co-authored-by: ThomasNing <thomas.ning@amd.com>

[ROCm/composable_kernel commit: c7b8e86e34]
2025-05-07 18:37:31 -07:00
Khushbu Agarwal
5d0b1b733c mfma_32x32x64_fp8/bf8 (#2148)
* support for mfma_32x32x64_fp8

* clang-formatted

* Fixing sparsity in codegen

[ROCm/composable_kernel commit: d58f2b8bd0]
2025-05-01 13:36:24 -07:00
Aviral Goel
b537ae86b9 [Tile Engine] Improved README.md (#2134)
* improved tile_engine readme

* changed ck tile explanation and json

* further improved readme

* fixed typo

[ROCm/composable_kernel commit: 1aea51d34e]
2025-04-29 17:37:07 -07:00
Khushbu Agarwal
aeb46e6a49 Support for MFMA_16x16x128 for fp8/bf8 (#2125)
* Adding 16x16x128 support for gfx950

* Support for fp8 and bf8

* fix input arguments for MFMA scale instruction

* clang-formatted

* Fixes for lwpck-3145 (#2138)

* Fix lds tile & cmake dep & default epilogue

* Fallback BTypeToUse to ADataType in WOQ cases

* reverting instance json file

* reverting instance json file

---------

Co-authored-by: Yi DING <yi.ding@amd.com>

[ROCm/composable_kernel commit: d107f3c3a5]
2025-04-28 18:19:50 -07:00
Khushbu Agarwal
10188b5103 [TileEngine] Support for sparsity in codegen (#2128)
* Added sparsity flag in codegen

* remove comments

* clan formatted

* added sparsity as runtime argument

* updated README

* updated stream config variable

* fix typo for tail_num in hot loop

[ROCm/composable_kernel commit: 768c99eca9]
2025-04-28 18:19:23 -07:00
Khushbu Agarwal
c6b7f48326 Adding include directory in tile_engine (#2116)
[ROCm/composable_kernel commit: 94662b02d0]
2025-04-22 15:55:19 -07:00
Khushbu Agarwal
74210a9dfc multi instance generation for CkTileEngine (#2080)
* Add support for multi-instance verification, print detail for each instance, documentation fix

* clang formatted

* Added Readme file

* updated readme

* Addressing review comments

* clang formatted

* Updated ReadMe and GPU reference code

* simplified dispatch kernel code

* indentation

[ROCm/composable_kernel commit: 7cadf187e2]
2025-04-21 08:39:45 -07:00
Khushbu Agarwal
09792fa112 file clang formatted (#2053)
[ROCm/composable_kernel commit: 3bda57c204]
2025-04-03 16:55:49 -07:00
Khushbu Agarwal
844730776f Documentation for newly added struct (#2051)
[ROCm/composable_kernel commit: b443056a26]
2025-04-03 16:24:34 -07:00
Khushbu Agarwal
b85b103194 [New] Build up the feature of CK Tile GEMM CodeGen (#1994)
* New branch for codegen changes

* Fix verify function for int4

* pk_int4 codegen

* Update to review comments

* Remove codegen directory and rename filenames

* Remove extra files; clean up CMake file

* New branch for codegen changes

* Fix verify function for int4

* pk_int4 codegen

* Update to review comments

* Remove codegen directory and rename filenames

* Remove extra files; clean up CMake file

* code changes for single instance

* config file rename, added few more combinations in json file

* Fix cmake file

* Addressing review comments

* Reverting files changed by merge to develop

---------

Co-authored-by: ThomasNing <thomas.ning@amd.com>

[ROCm/composable_kernel commit: fed0709121]
2025-04-03 11:54:12 -07:00