ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-11 17:00:18 +00:00
Code Issues Packages Projects Releases Wiki Activity
1,189 Commits 739 Branches 38 Tags
363feb482d03fa217adce6a56f5e74a2ae4a00c3
Commit Graph

1189 Commits

This Branch
This Branch
All Branches
Author SHA1 Message Date
Rostyslav Geyyer
363feb482d Refactor tolerances for correctness check in gemm op (#1188) 
* Refactor tolerances for correctness check

* Update tolerances

* Update host-side gemm

* Update reference gemm call
 2024-03-08 12:05:05 -08:00
Lisa
0e28de9766 Update link (#1186) 2024-03-07 10:09:17 -08:00
yhuiYH
adb3615d1a Update CODEOWNERS to use documentation group (#1190) 
Also had to remove a name
 2024-03-07 10:08:37 -08:00
dependabot[bot]
1ddc8a841a Bump rocm-docs-core from 0.35.0 to 0.35.1 in /docs/sphinx (#1187) 
Bumps [rocm-docs-core](https://github.com/RadeonOpenCompute/rocm-docs-core) from 0.35.0 to 0.35.1.
- [Release notes](https://github.com/RadeonOpenCompute/rocm-docs-core/releases)
- [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md)
- [Commits](https://github.com/RadeonOpenCompute/rocm-docs-core/compare/v0.35.0...v0.35.1)

---
updated-dependencies:
- dependency-name: rocm-docs-core
  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>
 2024-03-05 21:55:01 -08:00
Paul Fultz II
8eff4d62b6 Add host lib (#1134) 
* Format

* Format

* Format

* Remove const

* Use the right template

* Format

* Format

* add row/col instances

* Add missing file

* fixed

* Format

* Updates

* Format

* fixed rrr layout

* Format

* Update test and embed modules

* Restore older version

* Update year

* Set -fPIC

* Format

* Use double for isnan

* rename host folder to codegen + minor fix

* add codegen CI test

* add option to build components without building CK

* fix the groovy syntax

* fix typo

* use the correct function for the codegen stage

---------

Co-authored-by: Jing Zhang <jizha@amd.com>
Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
Co-authored-by: illsilin <Illia.Silin@amd.com>
 2024-03-05 17:08:43 -08:00
Illia Silin
cf86621170 [CI] Add CI build and test stage on MI300. (#1185) 2024-03-05 10:42:16 -08:00
Rostyslav Geyyer
9ce18b045d Fix example_gemm_xdl_fp8 (#1183) 2024-03-01 16:42:15 -08:00
Rostyslav Geyyer
acfb339238 Update clipping for fp8/bf8 conversion (#1182) 
* Update clipping for fp8 conversion

* Add clipping for bf8 conversion

* Format
 2024-03-01 10:30:38 -08:00
amoskvic
a776978cbe Style improvement: improving type alias usage consistency in gemm-related client examples. Also copyright year update for all client examples. (#1180) 
Co-authored-by: Arseny Moskvichev <amoskvic@amd.com>
 2024-02-28 16:39:03 -08:00
Illia Silin
d0c7b45150 Clip fp8 to +/-240 on all targets. (#1172) 
* clip fp8 to +/-240 on all targets

* if inputs to fp8 conversion are +/-inf, they remain unaltered

* increase tolerance for test_elementwise_layernorm to prevent false errors

* change the input values for gemm examples to floats

* reduce gemm example float input values to prevent errors

* increase the tolerance for gemm examples
 2024-02-27 14:31:05 -06:00
Bartłomiej Kocot
d909599729 Remove unnecessary comments (#1177) 2024-02-26 07:56:06 -06:00
dependabot[bot]
2eb74a9c0c Bump rocm-docs-core from 0.34.2 to 0.35.0 in /docs/sphinx (#1175) 
Bumps [rocm-docs-core](https://github.com/RadeonOpenCompute/rocm-docs-core) from 0.34.2 to 0.35.0.
- [Release notes](https://github.com/RadeonOpenCompute/rocm-docs-core/releases)
- [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md)
- [Commits](https://github.com/RadeonOpenCompute/rocm-docs-core/compare/v0.34.2...v0.35.0)

---
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>
 2024-02-22 21:39:16 -08:00
jakpiase
32d4be3d09 Add support for mixed precision bf16&int8 grouped gemm (#1166) 
* add support for mixed precision bf16&int8 grouped gemm

* fix gfx versions and add bf16 kbatch condition

* added reviewers comments
 2024-02-21 10:35:35 +01:00
Bartłomiej Kocot
66736edb95 Extend permute scale support up to 6D (#1168) 
* Extend permute scale support up to 6D

* Fixes

* Fixes

* Update profiler/README.md

Co-authored-by: Lisa <lisajdelaney@gmail.com>

* Update profiler/README.md

Co-authored-by: Lisa <lisajdelaney@gmail.com>

* Update profiler/README.md

Co-authored-by: Lisa <lisajdelaney@gmail.com>

* Update profiler/README.md

Co-authored-by: Lisa <lisajdelaney@gmail.com>

* Update profiler/README.md

Co-authored-by: Lisa <lisajdelaney@gmail.com>

* Update profiler/README.md

Co-authored-by: Lisa <lisajdelaney@gmail.com>

* Update profiler/README.md

Co-authored-by: Lisa <lisajdelaney@gmail.com>

---------

Co-authored-by: Lisa <lisajdelaney@gmail.com>
 2024-02-20 09:56:54 -08:00
dependabot[bot]
abac8b07dd Bump rocm-docs-core from 0.34.0 to 0.34.2 in /docs/sphinx (#1170) 
Bumps [rocm-docs-core](https://github.com/RadeonOpenCompute/rocm-docs-core) from 0.34.0 to 0.34.2.
- [Release notes](https://github.com/RadeonOpenCompute/rocm-docs-core/releases)
- [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md)
- [Commits](https://github.com/RadeonOpenCompute/rocm-docs-core/compare/v0.34.0...v0.34.2)

---
updated-dependencies:
- dependency-name: rocm-docs-core
  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>
 2024-02-16 07:48:52 -08:00
Illia Silin
b9ab9f4b4f upgrade the ccache version and update links (#1169) 2024-02-15 15:46:01 -08:00
Bartłomiej Kocot
1e73adbc28 Add optimized blockwise gemm using ck wrapper (#1157) 
* Add optimized blockwise gemm using ck wrapper

* Add basic gemm example

* Update docs

* Add tutorial for gemm using ck wrapper

* Add perf note

* edits

* Fix cmake

* Fixes

---------

Co-authored-by: Lisa Delaney <lisa.delaney@amd.com>
 2024-02-13 17:04:36 +01:00
Bartłomiej Kocot
bf98b47697 Add bilinear conv fwd and bwd data instances (#1164) 2024-02-13 11:49:05 +01:00
Illia Silin
a78be3f69e add docker credentials before pushing image (#1165) 2024-02-12 16:11:32 -08:00
zjing14
602c4cc0d9 Optimizing fp8_fp16 mixedprec gemm (#1150) 
* add delayed cvt

* extend fp16 gemm_splitk instances for fp8_fp16 gemm

* add f8 example

* add 128 kperblk instances for fp8

* add kpb128 instance

* added more instances into kpb128

* clean code

* clean code

* fix

* fix

* fixed

* Update example/35_splitK_gemm/splitK_gemm_xdl_fp16_fp8.cpp

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

* Update include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp

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

* Update library/src/tensor_operation_instance/gpu/gemm_splitk/device_gemm_xdl_splitk_f16_fp8_f16_mk_nk_mn_kpb128_instance.cpp

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

---------

Co-authored-by: Jing Zhang <jizha@amd.com>
Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>
 2024-02-12 09:45:42 -08:00
zjing14
94fbaac002 add generic instances for DeviceGemm_Xdl_CShuffle (#1161) 
* add generic instances

* clean code
 2024-02-09 10:20:53 -06:00
dependabot[bot]
4581b5d504 Bump rocm-docs-core from 0.33.2 to 0.34.0 in /docs/sphinx (#1163) 
Bumps [rocm-docs-core](https://github.com/RadeonOpenCompute/rocm-docs-core) from 0.33.2 to 0.34.0.
- [Release notes](https://github.com/RadeonOpenCompute/rocm-docs-core/releases)
- [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md)
- [Commits](https://github.com/RadeonOpenCompute/rocm-docs-core/compare/v0.33.2...v0.34.0)

---
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>
 2024-02-08 21:35:39 -08:00
Lakhinder Walia
1f306024d0 fast_gelu: minor code reorg to enhance ref & gpu performance (#1162) 2024-02-07 19:24:51 -08:00
Illia Silin
1b0fbaebbb Split-up instances to improve build times. (#1159) 
* split up splitk-gemm instances

* clean up some unused variables

* split the mk_kn_mn interwave splitk-gemm instances

* split up f16_f16_f16 mk_nk_mn splitk gemm instances

* fix clang format

* fix function names

* fix typo

* split up the 2 largest fp16*fp8 splitk gemm instances

* get rid of unused variables

* split up the largest splitk-gemm fp8*fp16 instance file

* split up the instances for xdl fp8 gemms

* split the headers for f16 and i8 for wmmma convolution instances
 2024-02-07 12:47:12 -08:00
jakpiase
ba86eadce5 Add support for mixed-precision f16bf16_int8 gemm (#1127) 2024-02-07 15:54:13 +01:00
dependabot[bot]
753cef783f Bump rocm-docs-core from 0.33.1 to 0.33.2 in /docs/sphinx (#1160) 
Bumps [rocm-docs-core](https://github.com/RadeonOpenCompute/rocm-docs-core) from 0.33.1 to 0.33.2.
- [Release notes](https://github.com/RadeonOpenCompute/rocm-docs-core/releases)
- [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md)
- [Commits](https://github.com/RadeonOpenCompute/rocm-docs-core/compare/v0.33.1...v0.33.2)

---
updated-dependencies:
- dependency-name: rocm-docs-core
  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>
 2024-02-06 21:24:32 -08:00
Bartlomiej Wroblewski
6951858221 Implement direct loads split-K GEMM kernel (#1137) 
* WIP: Implement direct loads split-K GEMM kernel

* Clean the review

---------

Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>
Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>
 2024-02-07 01:08:34 +01:00
dependabot[bot]
6299621107 Bump rocm-docs-core from 0.33.0 to 0.33.1 in /docs/sphinx (#1158) 
Bumps [rocm-docs-core](https://github.com/RadeonOpenCompute/rocm-docs-core) from 0.33.0 to 0.33.1.
- [Release notes](https://github.com/RadeonOpenCompute/rocm-docs-core/releases)
- [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md)
- [Commits](https://github.com/RadeonOpenCompute/rocm-docs-core/compare/v0.33.0...v0.33.1)

---
updated-dependencies:
- dependency-name: rocm-docs-core
  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>
 2024-02-06 09:24:34 -08:00
Illia Silin
f0dd1da088 Delete any dangling images after building a new one. (#1155) 
* delete dangling docker images

* fix groovy syntax

* fix groovy syntax again

* try a different way to delete dangling images
 2024-02-05 10:34:47 -08:00
Illia Silin
180f16f9ac Add support for more Navi2x and Navi3x models. (#1152) 
* add support for navi2x and navi3x models

* fix syntax

* use common macro for different mi300 architectures
 2024-02-02 11:35:26 -08:00
Bartłomiej Kocot
171ca260b5 Extend gemm traits number for ck wrapper (#1153) 2024-02-02 11:25:54 -08:00
Illia Silin
112b691bb7 add new performance tests for mixed fp16/fp8 gemms (#1151) 2024-01-31 13:27:17 -08:00
Bartłomiej Kocot
f3b6c23ac5 Add blockwise gemm to ck wrapper (#1139) 
* Add blockwise gemm to ck wrapper

* Add blockwise gemm traits

* Disable test_gemm for non xdl devices

* Fixes

* Add c layout descritpions
 2024-01-31 21:24:40 +01:00
Illia Silin
6651a124cc update the name of the compiler staging branch (#1148) 2024-01-30 13:55:31 -08:00
Illia Silin
e7495e6bb7 turn off performance tests in CI by default until the infrastructure is fixed (#1147) 2024-01-30 13:14:58 -08:00
dependabot[bot]
84832fc42d Bump rocm-docs-core from 0.31.0 to 0.33.0 in /docs/sphinx (#1144) 
Bumps [rocm-docs-core](https://github.com/RadeonOpenCompute/rocm-docs-core) from 0.31.0 to 0.33.0.
- [Release notes](https://github.com/RadeonOpenCompute/rocm-docs-core/releases)
- [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md)
- [Commits](https://github.com/RadeonOpenCompute/rocm-docs-core/compare/v0.31.0...v0.33.0)

---
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>
 2024-01-29 09:02:52 -08:00
Illia Silin
4a8297c28a fix CK path for hipTensor (#1143) 2024-01-25 17:05:43 -08:00
rocking
28f68a5a99 layernorm & groupnorm bwd gamma beta (#1133) 
* Add layernorm bwd gamma beta external api

* Add groupnorm external api

* Add layernorm bwd gamma beta profiler

* Add groupnorm bwd gamma beta ckProfiler

* Add layernorm & groupnorm bwd gamma beta test

* Fix groupnorm bwd gamma beta profiler bug

* Layernorm bwd weight client example

* Groupnorm bwd weight client example

* clang format

* Remove useless header

* Let inv_std be positive

* Rename to num_bytes and move this calculation outside the loop
 2024-01-25 19:53:15 +08:00
Illia Silin
180e572076 Fixing most of the cppcheck errors. (#1142) 
* fix cppcheck errors, first pass

* fix format

* fix returned value in examples

* add macro definitions for cppcheck

* fix the profile_gemm logic

* update the gemm profiler logic

* add more difinitions to cppcheck, fix couple more errors

* replace runtime error with message in device function

* fix a couple of int4 issues

* no return for fill function

* fix errors in data_types.hpp

* fix format

* fix few remaining errors

* fix errors in data_types.hpp

* fix last couple of errors in datat_types.hpp
 2024-01-24 13:47:48 -08:00
Bartłomiej Kocot
6169fbbdb3 Fix possible linting errors in changelog (#1141) 
* Fix possible linting errors in changelog

* Update CHANGELOG.md

* Update CHANGELOG.md

* Update CHANGELOG.md
 2024-01-24 17:19:02 +01:00
zjing14
1be4706366 fixed return (#1138) 2024-01-22 08:42:26 -08:00
Haocong WANG
bb63b9732c [GEMM] Optimization for MI200/300. (#1135) 
* Optimize GEMM on MI200/300:
1. Add new blockwise gemm pipeline
2. Add irregular splitk intances

* clang format + typo fix

* Fix a bug
 2024-01-19 07:02:22 -06:00
Bartłomiej Kocot
7e4eb4b800 Add optimized copy to ck wrapper (#1126) 
* Add optimized copy to ck wrapper

* Example optimizations

* Fixes

* Move img2col test to client example

* Refactor example

* Fix docs

* Fixes

* Fix

* Fixes

* Fixes

* Fixes

* Fixes

* Fixes

---------

Co-authored-by: zjing14 <zhangjing14@gmail.com>
 2024-01-19 11:29:00 +01:00
Illia Silin
38882d8ab5 add Adam to code owners (#1136) 2024-01-18 19:20:40 -06:00
randyh62
402a930a4a Randyh docfix (#1130) 
* Update LICENSE

update to 2024

* Update index.rst

change license.md to license.html

* fix syntax

---------

Co-authored-by: illsilin <Illia.Silin@amd.com>
 2024-01-16 09:00:37 -08:00
Illia Silin
c1b5b58192 add code owners (#1132) 2024-01-16 07:55:18 -08:00
Illia Silin
e6d099c830 Add cppcheck to CK CI. (#1125) 
* add cppcheck to the CK CI

* fix the path to CK source for cppcheck

* fix the path to CK source for cppcheck one more time

* fix the path to CK source for cppcheck third time

* change the path to ck_cppcheck.log

* install latest cppcheck from source

* fix bug in ck.hpp and use 20 threads for cppcheck

* create a switch to turn cppckeck on and off in CI
 2024-01-15 09:11:45 -08:00
dependabot[bot]
636a31015a Bump rocm-docs-core from 0.30.3 to 0.31.0 in /docs/sphinx (#1131) 
Bumps [rocm-docs-core](https://github.com/RadeonOpenCompute/rocm-docs-core) from 0.30.3 to 0.31.0.
- [Release notes](https://github.com/RadeonOpenCompute/rocm-docs-core/releases)
- [Changelog](https://github.com/RadeonOpenCompute/rocm-docs-core/blob/develop/CHANGELOG.md)
- [Commits](https://github.com/RadeonOpenCompute/rocm-docs-core/compare/v0.30.3...v0.31.0)

---
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>
 2024-01-15 09:09:13 -07:00
dependabot[bot]
0ce417269d Bump sphinxcontrib-bibtex from 2.6.1 to 2.6.2 in /docs/sphinx (#1129) 
Bumps [sphinxcontrib-bibtex](https://github.com/mcmtroffaes/sphinxcontrib-bibtex) from 2.6.1 to 2.6.2.
- [Changelog](https://github.com/mcmtroffaes/sphinxcontrib-bibtex/blob/develop/CHANGELOG.rst)
- [Commits](https://github.com/mcmtroffaes/sphinxcontrib-bibtex/compare/2.6.1...2.6.2)

---
updated-dependencies:
- dependency-name: sphinxcontrib-bibtex
  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>
 2024-01-11 11:27:03 -07:00
Illia Silin
886d9eeb99 Add an option to change the number of warm-up cycles and iterations. (#1124) 
* allow setting the number of warmup cycles and iterations for profiler

* fix the gemm_splitk and grouped_gemm examples
 2024-01-09 09:43:08 -08:00