Commit Graph

1779 Commits

Author SHA1 Message Date
Bartłomiej Kocot
b417137276 Grouped conv fwd v3 fix for SplitN an G > 1 (#2038)
* Grouped conv fwd v3 fix for SplitN an G > 1

* Remove int8 large test

* Retore int8 test

[ROCm/composable_kernel commit: ec742908bd]
2025-04-01 13:19:35 -07:00
Seunghoon Lee
0ac4c37028 Fix Windows build. (#2012)
* Remove duplicate using uint64_t.

* Cast before shift.

[ROCm/composable_kernel commit: df32020f93]
2025-04-01 12:22:10 -07:00
Max Podkorytov
70ad8571c2 add a fast compilation path for static for (0..N) (#2005)
* add a fast compilation path for static for (0..N)

* Update functional2.hpp

add comment and put range applier into detail namespace

* Update functional.hpp

ditto for ck-tile

* prettify

* prettify more

* add comment

* clang-format

[ROCm/composable_kernel commit: c59a8bb206]
2025-04-01 12:06:25 -07:00
Bartłomiej Kocot
a9ec282dba Improve compilation time for grouped conv fwd (#2039)
* Improve compilation time for grouped conv fwd

* Fix

[ROCm/composable_kernel commit: 6355ee7ca5]
2025-04-01 07:11:42 -07:00
Muhammed Emin Ozturk
30e5c8cb49 f8/bf16 GEMM Stream-K (#1879)
[ROCm/composable_kernel commit: dd4c12b155]
2025-03-31 20:30:17 -06:00
jefyang1
92f6e02b96 Fix gemm universal and grouped_conv_fwd test failures on gfx950 (#2031)
[ROCm/composable_kernel commit: 16b15e336a]
2025-03-31 09:20:52 -07:00
Adel Johar
de769ac359 Docs: Add precision support reference page (#1973)
* Docs: Add precision support reference page

* edit of the precision type content

* added more description on scalars

---------

Co-authored-by: spolifroni-amd <sandra.polifroni@amd.com>
Co-authored-by: Aviral Goel <aviral.goel@amd.com>

[ROCm/composable_kernel commit: fc073b483e]
2025-03-28 08:12:27 -06:00
rocking
2b657d9a2c Reduce redundant space in bias tensor (#2024)
Co-authored-by: Po Yen Chen <PoYen.Chen@amd.com>

[ROCm/composable_kernel commit: 8a20b62e91]
2025-03-28 21:58:06 +08:00
felix
817752cdb4 hotfix fix sorting int64 (#2025)
* fix sorting int64

* clang format

* fix example issue

* update WA issue #

---------

Co-authored-by: coderfeli <coderfeli@163.com>
Co-authored-by: carlushuang <carlus.huang@amd.com>

[ROCm/composable_kernel commit: a82f338fb9]
2025-03-28 11:31:52 +08:00
Illia Silin
5bfbcb38ac add gfx950 to default targets for rocm6.4+ (#2032)
[ROCm/composable_kernel commit: d142e15f5e]
2025-03-27 18:48:47 -07:00
spolifroni-amd
98cc377db2 creation of install doc and refactor of doc in general (#1908)
* creation of install doc and refactor of doc in general

* updates based on review comments

* updated based on review comments

* updated readme and contributors markdown

* added extra note to not use -j on its own

* added note about smoke tests and regression tests

* made changes as per Illia's feedback

---------

Co-authored-by: Aviral Goel <aviral.goel@amd.com>

[ROCm/composable_kernel commit: a426f67301]
2025-03-27 15:13:18 -06:00
felix
00342e69a2 ckmoe: change cmake; use smaller shape for i4 (#2027)
* change cmake; use smaller shape for i4

* fix pki4 run

* fix typo

* fix runtime arch logic for moe_gemm2 example

---------

Co-authored-by: coderfeli <coderfeli@163.com>
Co-authored-by: illsilin <Illia.Silin@amd.com>

[ROCm/composable_kernel commit: 36d50de50e]
2025-03-27 09:04:31 -07:00
Rostyslav Geyyer
48fa126a9e Add MX FP4 device conversion tests (#1889)
* Add conversion tests

* Fix ctor

* Fix nan logic

* Fix conversion logic

* Permute packed f4_t values

* Fix conversion to float, repack vector elements

* Fix device tests

* Permute elements in a vector

* Add a repro test

* Add a conversion for a repro test

* Update test vectors

* Update conversion

* Fix the test

* Update test vector generator

* Fix vector sr conversion

* Permute conversion args

* Update conversion

* Test

* Fix packing

* Simplify conversion function

* Pack conversion in a loop

* Pack conversion in a loop

* Pack another conversion in a loop

* Pack one more conversion in a loop

* Pack the last conversion in a loop

* Clean up

* Add printf to fix intrinsic

* Add a sw-based workaround

[ROCm/composable_kernel commit: 441343a23d]
2025-03-26 19:23:01 -05:00
Illia Silin
d424bbe440 Disable all pk_i4 tests for all targets except gfx942/950. (#2022)
* only build gemm_fp8_pk_i4 examples for gfx942/950

* fix cmake logic

* moved the architecture check to IsSupported function

* Revert "moved the architecture check to IsSupported function"

This reverts commit 056d2a08b3.

* disable all pk_i4 tests for targets other than gfx942/950

* fix cmake logic

[ROCm/composable_kernel commit: 23a949706c]
2025-03-26 15:15:57 -07:00
Bartłomiej Kocot
f967fd7296 Add support for GKCYX grouped conv fwd (#2015)
* Add support for GKCYX grouped conv fwd

* fixes

* fix

* changelog

* Fixes

[ROCm/composable_kernel commit: 54c81a1fcf]
2025-03-26 21:13:38 +01:00
Illia Silin
ba16351a03 fix clang format (#2021)
[ROCm/composable_kernel commit: fd915b83f7]
2025-03-26 09:42:10 -07:00
Mirza Halilčević
97b32a1f18 Add default arguments for prologue and epilogue. (#2020)
[ROCm/composable_kernel commit: 21e0ca197d]
2025-03-26 09:28:40 -07:00
Illia Silin
9a88442eba Make sure gemm_fp8_pk_i4 examples only build and run on gfx942/950. (#2010)
* only build gemm_fp8_pk_i4 examples for gfx942/950

* fix cmake logic

* moved the architecture check to IsSupported function

* Revert "moved the architecture check to IsSupported function"

This reverts commit 056d2a08b3.

[ROCm/composable_kernel commit: 99b2bbc1d6]
2025-03-25 14:43:38 -07:00
Andriy Roshchenko
3f06d019ba MX GEMM examples with FP8, FP16, and E8M0 scales (#2016)
* Add `scalar_type` specification for E8M0 exponent

* Specialize `nnvb_data_t_selector` for E8M0 exponent

* Remove partial specializations for `scalar_type` of `non_native_vector_base` template

* Reword command line helper string

* Create MX GEMM examples for different scales


[ROCm/composable_kernel commit: 72d888821c]
2025-03-25 15:33:03 -06:00
Illia Silin
aaf6a0343d Enable ClangBuildAnalizer when doing ninja build traces. (#2009)
* enable ClangBuildAnalizer when doing ninja traces

* add branch and date to clang build log name

* fix jenkins syntax

* fix jenkins syntax once more

* fix jenkins syntax once more

* simplify the clang_build log name

* simplify the clang_build log name further

[ROCm/composable_kernel commit: 44c093ba0c]
2025-03-25 12:27:04 -07:00
Max Podkorytov
a70c642a78 use fast path for sequence generation in old CK (#1993)
[ROCm/composable_kernel commit: 1a58522f01]
2025-03-25 11:28:44 -07:00
ruanjm
fcbf9630fe [CK_TILE] Improve RMS/Layer Normalization 2 Pass Pipeline Performance (#1861)
* 50ms -> 28ms

* Fix bug in non fuse_add_store cases

* Fine tuned setting for 2 pass pipeline

* adjust workload

* remove unnecessary change

* add layernorm

* Adding output quant and unquant results at the same time.

* fix test

* fix format

* tune for cases 128x640 and 128x1024

* bug ifx

[ROCm/composable_kernel commit: d49abdaa87]
2025-03-25 20:09:45 +08:00
Illia Silin
9dcfd79dcb Split up data_type header. (#1996)
* split fp64 vector data type

* add missing header

* move e8m0 structs

* split off numeric_utils header

* fix typo

* split off numeric limits header

* update data_type header

* fix clang format

* split off vector type header

* fix clang format

* fix typo for binary_inf

[ROCm/composable_kernel commit: d2eab23958]
2025-03-24 15:08:54 -07:00
Andriy Roshchenko
6b2b228eb4 Introduce MX GEMM for FP8 data type (#2000)
[ROCm/composable_kernel commit: 6660dc6b8e]
2025-03-24 15:41:07 -06:00
MHYang-gh
1998fcda26 Fix A/B lds transform (#2007)
[ROCm/composable_kernel commit: c027637a8f]
2025-03-22 23:13:50 -07:00
Bartłomiej Kocot
61b5e8318a Fix split N for large images in groupd conv fwd (#2004)
* Fix split N for large images in groupd conv fwd

* Fix comments

[ROCm/composable_kernel commit: 5b0873c31a]
2025-03-22 23:19:49 +01:00
carlushuang
20f331e7fe add mask support in hdim=192/128 (#1999)
[ROCm/composable_kernel commit: 6c08c5c46d]
2025-03-21 18:28:43 +08:00
BingYuan.Zhou
764d9468b2 fix ck_tile/basic_gemm build error (#1988)
[ROCm/composable_kernel commit: 5a0d693b86]
2025-03-20 22:01:14 -07:00
felix
6be18153f8 change cmake (#2006)
Co-authored-by: coderfeli <coderfeli@163.com>

[ROCm/composable_kernel commit: 902dbe89ad]
2025-03-20 19:25:11 -07:00
Attila T. Áfra
8cf49a44a9 Fix compile errors on Windows and Linux (#2002)
* Fix compile error on Windows (call to 'amd_wave_read_first_lane' is ambiguous)

* Fix compile error (no matching function for call to 'cast_to_f32_from_f8')

[ROCm/composable_kernel commit: c79bf11148]
2025-03-20 12:37:25 -07:00
carlushuang
f588e4b08e [CK_TILE] return value with macro in ck_tile::kernel_launch API (#1982)
* return value with macro and revert the return value

* [CK-TILE] no-macro launch api solution (#1992)

* no-macro solution

* address -Wcomma

---------

Co-authored-by: Max Podkorytov <4273004+tenpercent@users.noreply.github.com>

[ROCm/composable_kernel commit: e3c9886cdf]
2025-03-20 11:00:29 -07:00
jakpiase
9bf1c41338 [CK_TILE] Switch to universal gemm for batched and grouped gemms (#1919)
* switch to universal gemm for batched and grouped gemms

* added reviewer comments

* fixed grouped gemm tests

[ROCm/composable_kernel commit: 0e91d32c61]
2025-03-20 11:17:04 +01:00
rocking
8be61cfc9d Sync the kname with instance name (#1989)
Co-authored-by: Po Yen Chen <PoYen.Chen@amd.com>

[ROCm/composable_kernel commit: b819c217e4]
2025-03-20 00:06:45 +08:00
felix
7e4c8a56ed Ck moe hot fix (#1979)
* fix useless code and remove usless oob

* clang format

* fix coredump in e2e test

* fix2

* fix clang format

* fix output oob

* clang format

* rm useless comments

---------

Co-authored-by: coderfeli <coderfeli@163.com>
Co-authored-by: illsilin <Illia.Silin@amd.com>

[ROCm/composable_kernel commit: 7eaedeb36c]
2025-03-19 22:58:27 +08:00
Bartłomiej Kocot
71c6106947 Add grouped conv bwd wei merged grouped instance for larger filter (#1984)
* Add grouped conv bwd wei merged grouped instance for larger filter

* Update readme

[ROCm/composable_kernel commit: fdaff5603e]
2025-03-18 16:16:24 +01:00
Illia Silin
8df9260b3e Add a daily CI build on gfx908. (#1987)
* add one daily ci build on gfx908

* add redis invocation tag for gfx908

* make ci build for gfx908 conditional

* fix groovy logic

* add option to run perf tests for gfx908

* disable a few tests on mi100

[ROCm/composable_kernel commit: 1342ecf7fb]
2025-03-17 18:08:53 -07:00
Illia Silin
56ad0cdd62 disable ck_tile basic gemm (#1986)
[ROCm/composable_kernel commit: 07f25186b2]
2025-03-17 15:26:43 -07:00
aledudek
94a88e2ecc Async grouped gemm v3 (#1940)
* Fully async grouped gemm

* Remove commented code

* Remvoe maybe_unused

* host kernel args

* Checkpoint segfault debugging...

* Working part1

* Working part2

* Remvoe comments...

* Use void ptr for gemm kernel host args

* Fix device_grouped_gemm_multiple_d_dl build issue

* Fix device_grouped_gemm_xdl build issue

[ROCm/composable_kernel commit: 5095906975]
2025-03-17 16:42:43 +01:00
Bartłomiej Kocot
36f9cc5fb0 Grouped conv bwd data NGCHW (#1967)
* Grouped conv bwd data NGCHW

* fixes

* fix

* Improvements

* Fix

* Fix

* add client example

[ROCm/composable_kernel commit: c2e4898b4b]
2025-03-17 13:32:00 +01:00
valarLip
3d06952a2b hotfix fmoe build issue (#1976)
[ROCm/composable_kernel commit: 52b1cd7780]
2025-03-13 15:11:59 +08:00
dependabot[bot]
4643738f4d Bump rocm-docs-core from 1.17.1 to 1.18.1 in /docs/sphinx (#1977)
Bumps [rocm-docs-core](https://github.com/ROCm/rocm-docs-core) from 1.17.1 to 1.18.1.
- [Release notes](https://github.com/ROCm/rocm-docs-core/releases)
- [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md)
- [Commits](https://github.com/ROCm/rocm-docs-core/compare/v1.17.1...v1.18.1)

---
updated-dependencies:
- dependency-name: rocm-docs-core
  dependency-type: direct:production
  update-type: version-update:semver-minor
...

Signed-off-by: dependabot[bot] <support@github.com>
Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

[ROCm/composable_kernel commit: de7a745ca6]
2025-03-12 23:36:36 -07:00
carlushuang
5293517d0a Reapply "[CK_TILE] support hdim=192/128 pair for deepseekv3 (#1961)" … (#1971)
* Reapply "[CK_TILE] support hdim=192/128 pair for deepseekv3 (#1961)" (#1969)

This reverts commit a222573537b139c8a4f6b870910bdd487935efa8.

* fix codegen problem

* Update config.hpp

---------

Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>

[ROCm/composable_kernel commit: 3e81279d26]
2025-03-13 11:41:39 +08:00
Illia Silin
2559223c72 disable tests that take too long to build for gfx90a (#1975)
[ROCm/composable_kernel commit: d4a6d69643]
2025-03-12 17:54:03 -07:00
feli
4865ff49f3 ck_moe: fix useless code and remove usless oob (#1972)
* fix useless code and remove usless oob

* clang format

---------

Co-authored-by: coderfeli <coderfeli@163.com>

[ROCm/composable_kernel commit: 251afab3b7]
2025-03-12 09:22:42 -07:00
Illia Silin
73a1f03e58 use old instrinsics with staging compiler (#1970)
[ROCm/composable_kernel commit: 4c97cc511e]
2025-03-12 07:29:09 -07:00
Illia Silin
c62042b022 Revert "[CK_TILE] support hdim=192/128 pair for deepseekv3 (#1961)" (#1969)
This reverts commit 4399ad790293bb76e4cf89f333017cf06a6d3e82.

[ROCm/composable_kernel commit: 8cbcd3e0d0]
2025-03-11 10:40:18 -07:00
Haocong WANG
b31ac9d14d [Block Scale GEMM] Optimized block scale gemm (#1950)
* Added two kernel for M=32 problem

* Comment the first one

* Enable multiply_multiply for Scale_Block_M = 1 for deepseek

* Modify the a_thread offset since the A data load is different from B.

* edit fp8 ab scale for Scale_Block_M=1

* edit GemmSpec to MNKPadding

* enable blockwise pipelie v1 and v2. v1 is work for small K.

* add instance for gemm_ab_scale

* fix cmakelist of ckProfiler

* optimize blockscale gemm. todo: reduce vgpr usage

* fix a correctness bug

* sanity checked

* revert ckprofiler cmake changes

* clang format

* revert unnecessary changes.

* remove commented codes.

* split weight preshuffle library targets

* bring back enable-post-misched=0

* fix build issues for gemm_multiply_multiply_fp8 instances

* fix clang format

* add verbose build flag when building for all targets

* reduce path names for new instances

* fix paths in cmake

* refactor gemm_multiply_multiply library target

* fix a bug in example

* fix example 65 cmake

* reduce the number of threads when building libs for all targets to 50

* use ninja to build for all targets

* reduce teh number of threads when building for all targets

* reduce the number of threads to 32 when building libs for all targets to 50

---------

Co-authored-by: mtgu0705 <mtgu@amd.com>
Co-authored-by: chenjun <junchen2@amd.com>
Co-authored-by: illsilin <Illia.Silin@amd.com>
Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>

[ROCm/composable_kernel commit: cbd74c2d12]
2025-03-11 10:11:21 -07:00
Haocong WANG
6357ab039d reduce test size to avoid timeout on specific silicon (#1966)
[ROCm/composable_kernel commit: ba209b9dab]
2025-03-11 09:15:26 -07:00
Illia Silin
fa7f480e20 disable example_moe_gemm2_xdl_pk_i4 on gfx950 (#1968)
[ROCm/composable_kernel commit: aa42c3db06]
2025-03-11 08:34:47 -07:00
carlushuang
569e9892a6 [CK_TILE] support hdim=192/128 pair for deepseekv3 (#1961)
* support hdim=192/128 pair

* remove useless print

* update

[ROCm/composable_kernel commit: 7a93b16ff6]
2025-03-11 21:07:40 +08:00