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,825 Commits 776 Branches 40 Tags
moe_blockscale_func
Commit Graph

1825 Commits

This Branch
This Branch
All Branches
Author SHA1 Message Date
mtgu0705
f27fbdee49 fix the build flag for moe blockscale gfx950. 2025-05-03 08:48:58 -05:00
mtgu0705
6a0d86f683 function more reasonable. 2025-05-02 11:15:35 +08:00
mtgu0705
fe1648f0ce fix device kernel selection. 2025-04-30 13:33:31 +08:00
OscarXu
ac6672dd42 Merge remote-tracking branch 'refs/remotes/origin/moe_blockscale_dev' into moe_blockscale_dev 2025-04-29 16:11:41 +08:00
OscarXu
7024136ca1 Fix no-mainloop scenario 2025-04-29 16:09:49 +08:00
OscarXu
f3fedde7e6 Fix int overflow casued by large B matrix 2025-04-28 15:55:10 +08:00
OscarXu
481d158d61 Fix int overflow casued by large B matrix 2025-04-25 10:10:44 +08:00
OscarXu
e61179e734 Add v3 support. Add compile_option 2025-04-22 16:42:33 +08:00
mtgu0705
578945b346 fix a bug, function passed. 2025-04-22 14:35:53 +08:00
mtgu0705
d6b84c7b40 fix xdl transpose. 2025-04-22 09:46:37 +08:00
mtgu0705
2f6529dcc2 updated 2025-04-21 10:51:37 +08:00
mtgu0705
066f209640 added moe blockscale, compiler successful, function not correct. 2025-04-19 16:02:51 +08:00
mtgu0705
ef0b5f72af added moe_blockscale_dev branch to delete the per-token scale code. 2025-04-17 14:16:07 +08:00
mtgu0705
4dd4e8ed38 add example for moe stage2 fp8 blockscale 2025-04-14 10:50:28 +08:00
Illia Silin
3e6d21adeb enable gfx115x support (#2065) 2025-04-09 10:06:42 -07:00
MHYang-gh
03ce8729fd Make buffer coherence configurable in tensor view (#2041) 
* Make buffer coherence configurable in tensor view

* Fix clang-format for tensor_view.hpp
 2025-04-08 15:34:11 -07:00
valarLip
2c563fecf7 add passthrough for int32->float32 (#2062) 2025-04-08 15:16:30 -07:00
Khushbu Agarwal
263ff689e0 New instances for gemm_multiply_multiply_weightpreshuffle operator (#2061) 
* Add new instances for weight_preshuffle for f8->bf16

* Add new instances for weight_preshuffle for f8->f16

* clang formatted

---------

Co-authored-by: Khushbu Agarwal <khuagar@amd.com>
Co-authored-by: Thomas Ning <Thomas.Ning@amd.com>
 2025-04-08 15:14:53 -07:00
spolifroni-amd
2c8132126c fixed broken github link (#2063) 2025-04-08 10:20:31 -07:00
dependabot[bot]
b12cd6580b Bump rocm-docs-core from 1.18.1 to 1.18.2 in /docs/sphinx (#2047) 
Bumps [rocm-docs-core](https://github.com/ROCm/rocm-docs-core) from 1.18.1 to 1.18.2.
- [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.18.1...v1.18.2)

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

Signed-off-by: dependabot[bot] <support@github.com>
Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
 2025-04-08 09:06:38 -07:00
Max Podkorytov
6ce0797dad simplify generate_tuple (#2043) 2025-04-08 09:00:51 -07:00
aledudek
80aae6119b [CK_TILE] Fix GEMM Memory Pipeline (#2034) 
* [CK_TILE] Fix GEMM Memory Pipeline

* Fix transpose tile

* Add comments
 2025-04-08 12:40:04 +02:00
Illia Silin
72c0261ef1 Fix a couple of CI issues. (#2050) 
* fix jenkins jobs

* fix perf log name for gfx908

* only run gemm perf tests on gfx908
 2025-04-07 12:48:34 -07:00
Illia Silin
1793228422 fix codegen issues (#2052) 2025-04-07 07:08:39 -07:00
Illia Silin
29f7266216 Revert "CkProfiler StreamK GemmUniversal Fix and Split Gemm_universal Test (…" (#2054) 
This reverts commit 7142d8003c.
 2025-04-07 06:49:36 -07:00
slippedJim
5a22b61de5 Add new receipt (#2055) 2025-04-07 14:18:01 +08:00
Khushbu Agarwal
3bda57c204 file clang formatted (#2053) 2025-04-03 16:55:49 -07:00
Khushbu Agarwal
b443056a26 Documentation for newly added struct (#2051) 2025-04-03 16:24:34 -07:00
Illia Silin
572cd820ce Split env.hpp header from the ck.hpp header. (#2049) 
* split env.hpp out of main headers

* fix namespace logic
 2025-04-03 15:30:21 -07:00
Muhammed Emin Ozturk
7142d8003c CkProfiler StreamK GemmUniversal Fix and Split Gemm_universal Test (#2044) 
* fix and split gemm_universal test

* clang

* Update test_gemm_universal_ut_cases_bf16.inc

* Update test_gemm_universal_xdl_bf16.cpp

* Update test_gemm_universal_ut_cases_fp16.inc
 2025-04-03 14:22:43 -07:00
Khushbu Agarwal
fed0709121 [New] Build up the feature of CK Tile GEMM CodeGen (#1994) 
* New branch for codegen changes

* Fix verify function for int4

* pk_int4 codegen

* Update to review comments

* Remove codegen directory and rename filenames

* Remove extra files; clean up CMake file

* New branch for codegen changes

* Fix verify function for int4

* pk_int4 codegen

* Update to review comments

* Remove codegen directory and rename filenames

* Remove extra files; clean up CMake file

* code changes for single instance

* config file rename, added few more combinations in json file

* Fix cmake file

* Addressing review comments

* Reverting files changed by merge to develop

---------

Co-authored-by: ThomasNing <thomas.ning@amd.com>
 2025-04-03 11:54:12 -07:00
Thomas Ning
50d1f8ff90 Add the MI355 support for CK TILE GEMM (#2046) 
* Get the root cause of the ck tile gemm failing on mi355

* Fix the ck tile gemm on MI355

* delete the debug info
 2025-04-03 11:48:54 -07:00
Rostyslav Geyyer
265af71a71 Add FP16/BF16<->FP8/BF8 conversions (#2035) 
* Move conversion functions and add missing conversions

* Add tests

* Add missing conversions

* Add missing conversions

* Add bf8 tests

* Update clipping for vectors

* Add missing conversions

* Add bf16 fp8 tests

* Add bf16 bf8 tests

* Fix device conversion

* Fix conversions

* Fix vector use

* Minor fix

* Add a workaround flag

* Add a workaround flag for bf16 conversion

* Add another workaround

* Add a workaround for fp16 to bf8 conversion

* Update type alias

* Add docstrings and missing wrappers

* Fix if defined macros

* Fix more if defined macros

* Add comments

* Remove __host__ specifier

* Add a gfx950 guard

* Update function naming
 2025-04-03 12:42:03 -05:00
aledudek
9329432f6c Post-merge changes for fully async args copy in ck grouped gemm (#1991) 
* Post-merge changes for fully async args copy in ck grouped gemm

* Post-merge documentation and naming changes

* Build fix and updated changelog

* Revised comments
 2025-04-03 13:35:43 +02:00
Bartłomiej Kocot
2ccf914888 Add support for GKCYX grouped conv weight (#2023) 
* Grouped conv bwd weight GKCYX support

* fix and changelog

* fix

* fix

* fixes

* comments

* fix
 2025-04-02 23:59:49 +02:00
Adam Osewski
e5ad48a784 Basic docs for universal gemm & ck-tile gemm. (#2014) 
* Basic docs for universal gemm & ck-tile gemm.

* Update include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>

* Update include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>

* Update include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>

* Update include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>

* Update include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>

* Update include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>

* Update include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>

* Update include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>

* Update include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>

* Update include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>

* Update include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>

* Update include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>

* Update include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>

* Update include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>

* Update include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>

* Update include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>

* Update include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>

* Update include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>

* Update include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>

* Update include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>

* Update include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp

Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>

* Reviewers suggestions.

* Align tparam names in doc with class tparams.

* More reviewers fine tuning ;)

---------

Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>
Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>
 2025-04-02 11:03:40 +02:00
Bartłomiej Kocot
8c0ab61ece Grouped conv backward data GKCYX support (#2029) 
* Grouped conv backward data GKCYX support

* profiler

* Converter

* split instances
 2025-04-01 13:24:38 -07:00
Bartłomiej Kocot
ec742908bd 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
 2025-04-01 13:19:35 -07:00
Seunghoon Lee
df32020f93 Fix Windows build. (#2012) 
* Remove duplicate using uint64_t.

* Cast before shift.
 2025-04-01 12:22:10 -07:00
Max Podkorytov
c59a8bb206 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
 2025-04-01 12:06:25 -07:00
Bartłomiej Kocot
6355ee7ca5 Improve compilation time for grouped conv fwd (#2039) 
* Improve compilation time for grouped conv fwd

* Fix
 2025-04-01 07:11:42 -07:00
Muhammed Emin Ozturk
dd4c12b155 f8/bf16 GEMM Stream-K (#1879) 2025-03-31 20:30:17 -06:00
jefyang1
16b15e336a Fix gemm universal and grouped_conv_fwd test failures on gfx950 (#2031) 2025-03-31 09:20:52 -07:00
Adel Johar
fc073b483e 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>
 2025-03-28 08:12:27 -06:00
rocking
8a20b62e91 Reduce redundant space in bias tensor (#2024) 
Co-authored-by: Po Yen Chen <PoYen.Chen@amd.com>
 2025-03-28 21:58:06 +08:00
felix
a82f338fb9 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>
 2025-03-28 11:31:52 +08:00
Illia Silin
d142e15f5e add gfx950 to default targets for rocm6.4+ (#2032) 2025-03-27 18:48:47 -07:00
spolifroni-amd
a426f67301 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>
 2025-03-27 15:13:18 -06:00
felix
36d50de50e 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>
 2025-03-27 09:04:31 -07:00
Rostyslav Geyyer
441343a23d 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
 2025-03-26 19:23:01 -05:00