Qun Lin
3225f42dab
update script
2025-06-06 17:34:22 +08:00
Lin, Qun
badd82498a
fix bugs 2
2025-06-06 04:31:39 -05:00
Lin, Qun
9b7260f544
fix bugs in wsplit and k-split
2025-06-05 23:37:23 -05:00
Qun Lin
7f93429797
add wsplit
2025-06-05 17:40:48 +08:00
Qun Lin
fca03f72fc
support k-split
2025-06-05 13:05:49 +08:00
Lin, Qun
289fdf59e1
finalize code
2025-06-03 23:25:07 -05:00
Lin, Qun
a61d19ff8f
fix a typo
2025-06-03 19:28:10 -05:00
Qun Lin
9f45a878fc
add test script
2025-06-03 21:31:23 +08:00
Qun Lin
5d1c8c826f
add pipelinv 2
2025-06-03 21:30:22 +08:00
Lin, Qun
4c39de8b8e
fix bugs
2025-06-03 06:26:42 -05:00
Lin, Qun
3f18187dd9
fix bug 2
...
- kernel 1-3 passed
2025-06-03 02:21:31 -05:00
Lin, Qun
3516dbdf56
fix bugs.
2025-06-03 01:24:45 -05:00
Qun Lin
6d12067a30
fix a typo
2025-06-02 16:17:58 +08:00
Qun Lin
eb158d9b86
add device factory and fix build error.
2025-06-02 15:05:10 +08:00
Qun Lin
b27dddfb59
bwd weight dl v4 - v0
2025-06-02 10:39:13 +08:00
Qun Lin
044df51b56
add new class
2025-05-30 20:25:44 +08:00
Lin, Qun
75b6cd4934
fix quality issue
...
option MergeShareMem doesn't work when Num_Wave > 1
2025-05-29 23:28:37 -05:00
Lin, Qun
26b8a1d972
change layout to gnhwc
2025-05-29 19:28:38 -05:00
Lin, Qun
321dc5dd39
version 3
2025-05-29 08:25:18 -05:00
Lin, Qun
5fd4aef2dc
version 2
2025-05-29 07:38:04 -05:00
Lin, Qun
0750ebb13d
version 1
2025-05-29 01:08:04 -05:00
Qun Lin
b9466aff4d
tmp code 2
2025-05-29 09:01:03 +08:00
Qun Lin
45b4c48000
tmp code 1
2025-05-28 15:47:05 +08:00
Lin, Qun
2e57e6e61c
refine example_conv_bwd_weight
2025-05-27 08:54:35 +08:00
Lin, Qun
f5377edfe5
refine test example_grouped_conv_bwd_weight_xdl_fp16
2025-05-27 08:54:35 +08:00
Lin, Qun
94ae9ae73f
update ck_tile script
2025-05-27 08:54:33 +08:00
Qun Lin
2f0015e788
disable unsupported cases
2025-05-27 08:53:51 +08:00
Qun Lin
7e919d67f2
disable unuspport cases
2025-05-27 08:53:51 +08:00
Qun Lin
11daf454b9
disable unsupported cases.
2025-05-27 08:53:51 +08:00
Lin, Qun
fb212bd921
fix ck tile build error
2025-05-27 08:53:50 +08:00
Qun Lin
027d58110a
disable test to pass compile.
2025-05-27 08:52:30 +08:00
Qun Lin
acf28d2c41
rebase flatmm fp8
2025-05-27 08:52:28 +08:00
Qun Lin
feeb6f7cd8
refine universal gemm multi-config
2025-05-27 08:49:57 +08:00
Qun Lin
303336858f
fix test script2
2025-05-27 08:49:57 +08:00
Qun Lin
512e0d7935
fix test scripts
2025-05-27 08:49:57 +08:00
Qun Lin
b1d38e9929
add ck tile test script
2025-05-27 08:49:57 +08:00
Qun Lin
9c37603836
extend scrip to conver all ck profiler op
2025-05-27 08:49:57 +08:00
Qun Lin
547c602d85
support multi config in universal gemm
2025-05-27 08:49:57 +08:00
Lin, Qun
16bb17dc71
add config 7
2025-05-27 08:49:57 +08:00
Lin, Qun
ede9f6ca8b
add config 3.4.5.6
2025-05-27 08:49:57 +08:00
Qun Lin
efe964bd95
1. copy bhalf v3 to half v3
...
2. add multi config in half v2
3. sync block universal with half v3
2025-05-27 08:49:54 +08:00
Bartłomiej Kocot
b1ed92b131
Revert "Remove not needed bwd wei merged groups instances ( #2218 )" ( #2235 )
...
This reverts commit 4583aeffad .
2025-05-26 23:26:04 +02:00
Bartłomiej Kocot
4583aeffad
Remove not needed bwd wei merged groups instances ( #2218 )
...
* Grouped conv bwd wei add two stage instances for larger filter and Merge Groups
* Fix
* fix
* Revert "Restore oddc instances (#2201 )"
This reverts commit 6342f6b5e8 .
* fix
---------
Co-authored-by: Rostyslav Geyyer <46627076+geyyer@users.noreply.github.com >
2025-05-26 22:46:18 +02: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
Illia Silin
bc2551ac3b
disable building device_mha_operations by default ( #2225 )
2025-05-22 14:03:04 -07:00
Adam Dickin
417a6b65b6
Add MIOPEN_REQ_LIBS_ONLY option for cmake to build only the libs MIOpen requires ( #2224 )
...
* cut out anything we dont need for MIOpen to test
* refactor exclusion code to be more streamlined.
2025-05-22 11:14:33 -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