Commit Graph

538 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
Bartłomiej Kocot
037764bbc6 Fix grid size calc for bwd wei (#2226) 2025-05-26 16:51:09 +02: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
jefyang1
f18170064d Use new mfma instructions for FP8 on gfx950 (#2202)
* Add logic to use new mfma instructions for fp8 bf8

* Fix example_gemm_xdl_fp8_pk_i4_bpreshuffle_v3 on gfx950 and run clang format

* Update include/ck/tensor_operation/gpu/warp/xdlops_gemm.hpp

Co-authored-by: Andriy Roshchenko <107577548+andriy-ca@users.noreply.github.com>

* Fix intrin_mfma f8 calls due to merge mistake

---------

Co-authored-by: Andriy Roshchenko <107577548+andriy-ca@users.noreply.github.com>
2025-05-19 17:29:51 -07:00
arai713
5b3430b868 Narrowing error fix for codegen compilation (#2194)
* removed comment with special characters

* fix for arg/template change after merge from develop

---------

Co-authored-by: Thomas Ning <Thomas.Ning@amd.com>
2025-05-16 11:11:54 -07:00
Mateusz Ozga
fa3c6811d8 Disable conv for Filter1x1Stride1Pad0 when K or C is even (#2186) 2025-05-16 10:18:47 +02:00
Bartłomiej Kocot
c53b7bd22e Switch to v2 pipeline for grouped conv bwd data (#2181)
* Change to old pipeline for grouped conv bwd data

* fix

* fix

* fix

* fix

* fix

* fix

* Fix
2025-05-13 10:14:30 +02:00
Thomas Ning
b49f7de81f Improve the general performance of the Preshuffled GEMM V3 & delete the unnecessary instances (#2166)
* make the work compiled

* Solved the example code, but still have the profiler error

* Finished the feature

* Clang format and update the CHANGELOG

* solve the preshuffle v1 & v2 problem

* Comment Addressed

* Comment Addressed
2025-05-12 09:52:58 -07:00
Bartłomiej Kocot
6fddb5708c Add grouped conv fwd bias relu instances (#2179)
* Add grouped conv fwd bias relu instances

* fixes

* fix
2025-05-09 22:52:34 +02:00
jefyang1
6b1a339b6f Fix grouped conv bwd data tests on gfx950 (#2173) 2025-05-09 09:01:06 -07:00
Andriy Roshchenko
cb27e7c77f Ensure MX GEMM Instances can be Cross-Compiled for Multiple Architectures (#2171)
* Re-enable MX GEMM instances

* Fix compilation error when building MX GEMM for multiple architectures
2025-05-08 13:26:03 -06:00
Andriy Roshchenko
79b0bfeb41 MX GEMM - Add FP8 GEMM Tests for Different Layouts (#2152)
* Add gemm_mx_fp8_bf8 example with row-major B

* Add more overloads of MX MFMA instructions

* Add MK_KN (RRR) tests

* Add KM_NK (CCR) tests

* Add more problem sizes to Large tests

* Add test_gemm_mx to the list of regression tests
2025-05-01 11:55:48 -06:00
Bartłomiej Kocot
4094ad158a Integrate universal gemm with conv bwd data and add SplitK (#1315)
* Integrate universal gemm with conv bwd data

* Fix multi d kernel

* Add splitK support

* instances refactor

* instances refactor

* refactor

* fixeS

* fixes

* 16x16 instnaces

* Fixes

* Fix

* Fix

* Fix

* Fix

* Fix

* Fixes

* fix

* fix
2025-04-28 23:54:49 +02:00