Commit Graph

645 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
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
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
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
Andriy Roshchenko
57e0f5df29 MX GEMM - Expand MX MFMA Testing to BF8, FP6, and BF6 Data Types (#2199)
* Unify test interface for different layouts.

* WIP: Introducing FP4/FP6/FP8 abstractions

* WIP: Introducing packed storage abstraction

* WIP: Introducing packed storage abstraction

* WIP: Improved support for FP6 data type

* Refactor packed storage for f6_t

* WIP: FP6 MFMA test

* Test if we correctly represent all FP6/FP4 numbers

* Additional output for failed FP4 test.

* More failing conversion tests

* Even more failing conversion tests

* Working FP6 MFMA tests

* Expand MX MFMA testing to BF8/6

* Update and verify MX MFMA test for packed types

* Fix fp4 and fp6 conversions on host

* Working MX MFMA tests for FP8/6/4

* Cleanup

* Add missing type

* Cleanup

* Final cleanup

* Restrict FP6/4 values output to CK_LOGGING=1

* Use CHAR_BIT instead of number 8

* Fix typo

* Remove FP6 and FP4 from the list of native types

---------

Co-authored-by: Rostyslav Geyyer <rosty.geyyer@amd.com>
2025-05-19 16:52:51 -05: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
Illia Silin
58f9e9ffbc Update the buffer load/store intrinsic names for clang>=20. (#2192)
* fix the buffer load/store intrinsic names

* fix clang format
2025-05-13 10:18:14 -07: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