ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-29 11:16:59 +00:00
Code Issues Packages Projects Releases Wiki Activity
1,976 Commits 776 Branches 40 Tags
bwd_data_1c_dev
Commit Graph

841 Commits

Author SHA1 Message Date
joyeamd
c02516c8b8 fix a compiling issue 2025-06-13 13:26:48 +08:00
joyeamd
a7e874b13f add back original kernel 2025-06-13 13:21:09 +08:00
joye
8d6ca68d90 update kernel 2025-06-11 14:40:30 +08:00
joye
09a2e029b9 add lds barrier 2025-06-11 12:44:03 +08:00
joye
4c85e10281 update kernel 2025-06-11 12:38:22 +08:00
joye
7bc604f06a update kernel 2025-06-11 11:10:16 +08:00
joye
b90beae6f0 update kernel 2025-06-11 11:05:49 +08:00
joye
f0a0e79d67 fix compiling errors 2025-06-11 11:01:25 +08:00
joye
0eefe9ac1f update kernel 2025-06-11 10:55:28 +08:00
joye
376e0992ef fix compiling errors; now can pass compilation 2025-06-10 08:34:27 +08:00
joyeamd
4e469f5572 update kernel 2025-06-10 08:11:30 +08:00
joyeamd
d70acd683b update kernel 2025-06-09 21:13:02 +08:00
joyeamd
4c4dd342ed fix compiling errors 2025-06-09 20:44:03 +08:00
joyeamd
9872d2e159 fix some compiling errors 2025-06-09 20:36:55 +08:00
joyeamd
b64a2cfda8 fix some compiling errors 2025-06-09 20:29:30 +08:00
joyeamd
e43bc46629 update kernel 2025-06-09 20:24:47 +08:00
joyeamd
b1d03f7e8a update kernel 2025-06-09 19:37:44 +08:00
joye
0677989d23 Merge branch 'bwd_data_1c_dev' of github.com:ROCm/composable_kernel into bwd_data_1c_dev 2025-06-09 15:34:21 +08:00
joye
51dad1aaca update shader 2025-06-09 15:34:04 +08:00
joyeamd
9b34909b76 add default nullptr 2025-06-09 15:13:40 +08:00
joyeamd
c631001f4e update kernel 2025-06-09 14:51:31 +08:00
joye
48a0cee750 update kernel 2025-06-09 14:41:32 +08:00
joye
1e444a8b27 update kernel 2025-06-09 14:09:33 +08:00
joye
d19a6c0ef0 update kernel 2025-06-09 13:45:31 +08:00
joye
cdf21a7178 update kernel 2025-06-08 20:43:19 +08:00
joye
4fe245e8d5 update kernel; not correctly 2025-06-06 18:35:13 +08:00
joye
a3934c5141 update shader 2025-06-05 16:26:00 +08:00
joye
65df6b65ed update variable to not hard coding 2025-06-05 15:09:35 +08:00
joye
69b6a8b20c update shader 2025-06-04 15:18:00 +08:00
joye
075527783c update shader 2025-06-04 12:21:38 +08:00
joye
37555a8f66 the current best kernel 2025-06-03 15:56:33 +08:00
joye
850b9adbf9 current perf best kernel 2025-06-03 12:56:38 +08:00
joye
d50a7ac6cb update problem shader 2025-06-03 11:08:26 +08:00
joye
f8bc223a58 update kernel to pass 2025-06-03 09:27:57 +08:00
joye
8553458b9b fix some compiling errors 2025-06-03 09:00:47 +08:00
joye
2a3dcc6d51 delete a typo error 2025-06-03 08:54:34 +08:00
joye
945e3a44ad add another kernel 2025-06-03 08:37:20 +08:00
Bartłomiej Kocot
e7906dd644 Change relu to clamp for grouped conv fwd instances (#2249) 2025-05-29 00:51:25 +02:00
Illia Silin
bbdaf79a52 Revert "[CK_tile] Add rotating buffer feature for universal gemm (#2200)" (#2256) 
This reverts commit 99857e10e6.
 2025-05-28 09:46:52 -06:00
Khushbu Agarwal
99857e10e6 [CK_tile] Add rotating buffer feature for universal gemm (#2200) 
* Add rotating buffer feature for universal gemm

* adding changes in tile_engine

* Updated code to merge kernel_launch

* removing comments

* Enable rotating buffer changes to flatmm

* Created diff launch_kernel function for rotating buffer

* Simplfied calculation using macros

* merge code with new changes in tile_engine

* clang formatted

* Redefine macros
 2025-05-27 23:00:58 -07:00
Casey-Shi
128f5a1eab [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>
 2025-05-26 22:32:36 -07:00
Po Yen Chen
c42b957d65 [CK_TILE] For FMHA forward kernels, assign block indices reversely if using mask (#2209) 
* Assign block indices reversely if kHasMask=true

* Assign block indices reversely for splitkv kernel
 2025-05-27 10:58:58 +08:00
Bartłomiej Kocot
037764bbc6 Fix grid size calc for bwd wei (#2226) 2025-05-26 16:51:09 +02:00
Zzz9990
ece38b9d7a [VLLM V1] Add chunked prefill for FA to pass seq with small seqlen_q (#2221) 
* fix splitkv compiler issue since lse is used to select kernel instances

* bypass seqlen == 1

* add chunked prefill into mha varlen

This reverts commit aa9847e42d.

* skip compile when receipt 2-4 and add comments

* fix

---------

Co-authored-by: fsx950223 <fsx950223@outlook.com>
 2025-05-26 19:17:18 +08:00
Illia Silin
8146e471f1 fix the buffer intrinsic names for clang >=20 (#2228) 2025-05-23 14:58:25 -07:00
Illia Silin
1b846143c6 Revert "Update the buffer load/store intrinsic names for clang>=20. (#2192)" (#2227) 
This reverts commit 58f9e9ffbc.
 2025-05-22 15:41:17 -07:00
Aviral Goel
534d4594d0 Refactor tile_window.hpp, tile_window_linear.hpp into a CK Tile Hierarchy (#2214) 
* window_origin variable now in base class

* abstracted more functions

* consolidated tile_window_static_distribution and tile_window_static_lengths

* clang format

* skeleton code for tile_window and tile_window_linear consolidation

* more abstraction

* moved variables from child to parent

* clang format

* removed comments

* removed debug code

* removed debug code

* abstracting traits WIP

* consolidated traits

* removed comments and clang formatted
 2025-05-21 23:28:00 -07:00
Aviral Goel
fa39c4e798 Add Doxygen Documentation for HostTesnor, HostTensorDescriptor, DeviceMem, FillUniformDistribution (#2160) 
* added documentation for HostTensorDescriptor

* added documentation for DeviceMem and FillUniformDistribution

* fixed merging error

* fixed host_tensor_descriptor error

* clang format
 2025-05-21 10:34:30 -07:00
Thomas Ning
1386924749 Add the instances for small sized GEMM in preshuffle and improve CMake Flag (#2212) 
* Add small instance, add the bug fix, & improve the example CMake

* clang format
 2025-05-20 15:05:08 -07:00
Sami Remes
d1e6f0982d [CK_TILE] Grouped GEMM tile loop (#2146) 
* Add trait to use a persistent kernel and split the entrypoints in grouped gemm

* Some helper functions for persistent kernel case

* Get max occupancy grid using device properties

* Implement tile loop in main entry point to grouped gemm

* Enable GridSize() on device

* Handle offset tile index using real current block index

* Add persistent kernel choice to grouped gemm example

* Use a for-loop for iterating over the group

* Reduce VGPR spills by early-exit

* Enable persistent kernel choice in grouped_gemm example

* Add persistent kernel option to grouped_gemm test

* Fix formatting with remod.py

* Remove GridUpdateBlocks as blocks are now iteratively computed

* Add comment about VGPR spilling

* Fix formatting

* Use CK_TILE_HOST instead of __host__

* Enable all Row/Col combinations in grouped gemm unit test

* Add some KBatch=2 cases to grouped gemm tests

* Fix SplitK for grouped gemm

* Enable pipeline hotloop/tailnumber selection in-kernel for grouped gemm

* Add type traits

* Split examples to regular and tileloop

* Formatting

* Use hipExtStreamGetCUMask to get current active CUs for the given stream

* Align test and example kernel config, and disable validation for splitk repeats

* Remove debug options from CMakeLists.txt

* Separate the code paths for persistent/non-persistent in test

* Fix formatting

* Address review comments

---------

Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>
 2025-05-20 17:18:57 +03:00