ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-14 02:02:46 +00:00
Code Issues Packages Projects Releases Wiki Activity
1,174 Commits 742 Branches 38 Tags
8921adfb3934c5decd94f276c7f89c882563c019
Commit Graph

1174 Commits

This Branch
This Branch
All Branches
Author SHA1 Message Date
Illia Silin
8921adfb39 upgrade the ccache version and update links (#1169) 
[ROCm/composable_kernel commit: b9ab9f4b4f]
 2024-02-15 15:46:01 -08:00
Bartłomiej Kocot
22bfdc1168 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>

[ROCm/composable_kernel commit: 1e73adbc28]
 2024-02-13 17:04:36 +01:00
Bartłomiej Kocot
34bb3c7e31 Add bilinear conv fwd and bwd data instances (#1164) 
[ROCm/composable_kernel commit: bf98b47697]
 2024-02-13 11:49:05 +01:00
Illia Silin
baa25bcc4b add docker credentials before pushing image (#1165) 
[ROCm/composable_kernel commit: a78be3f69e]
 2024-02-12 16:11:32 -08:00
zjing14
8063a407e7 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>

[ROCm/composable_kernel commit: 602c4cc0d9]
 2024-02-12 09:45:42 -08:00
zjing14
cc25dc316f add generic instances for DeviceGemm_Xdl_CShuffle (#1161) 
* add generic instances

* clean code

[ROCm/composable_kernel commit: 94fbaac002]
 2024-02-09 10:20:53 -06:00
dependabot[bot]
ec7641b27b 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>

[ROCm/composable_kernel commit: 4581b5d504]
 2024-02-08 21:35:39 -08:00
Lakhinder Walia
77addd3131 fast_gelu: minor code reorg to enhance ref & gpu performance (#1162) 
[ROCm/composable_kernel commit: 1f306024d0]
 2024-02-07 19:24:51 -08:00
Illia Silin
0a90ec2237 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

[ROCm/composable_kernel commit: 1b0fbaebbb]
 2024-02-07 12:47:12 -08:00
jakpiase
b41cf51951 Add support for mixed-precision f16bf16_int8 gemm (#1127) 
[ROCm/composable_kernel commit: ba86eadce5]
 2024-02-07 15:54:13 +01:00
dependabot[bot]
f59cd842c8 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>

[ROCm/composable_kernel commit: 753cef783f]
 2024-02-06 21:24:32 -08:00
Bartlomiej Wroblewski
dd3c37d031 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>

[ROCm/composable_kernel commit: 6951858221]
 2024-02-07 01:08:34 +01:00
dependabot[bot]
e000e4c442 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>

[ROCm/composable_kernel commit: 6299621107]
 2024-02-06 09:24:34 -08:00
Illia Silin
6d7cc0a989 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

[ROCm/composable_kernel commit: f0dd1da088]
 2024-02-05 10:34:47 -08:00
Illia Silin
b63148b2b1 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

[ROCm/composable_kernel commit: 180f16f9ac]
 2024-02-02 11:35:26 -08:00
Bartłomiej Kocot
1971822447 Extend gemm traits number for ck wrapper (#1153) 
[ROCm/composable_kernel commit: 171ca260b5]
 2024-02-02 11:25:54 -08:00
Illia Silin
6b17c34c28 add new performance tests for mixed fp16/fp8 gemms (#1151) 
[ROCm/composable_kernel commit: 112b691bb7]
 2024-01-31 13:27:17 -08:00
Bartłomiej Kocot
240330481c 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

[ROCm/composable_kernel commit: f3b6c23ac5]
 2024-01-31 21:24:40 +01:00
Illia Silin
a2bc19ab69 update the name of the compiler staging branch (#1148) 
[ROCm/composable_kernel commit: 6651a124cc]
 2024-01-30 13:55:31 -08:00
Illia Silin
94cb9d6d85 turn off performance tests in CI by default until the infrastructure is fixed (#1147) 
[ROCm/composable_kernel commit: e7495e6bb7]
 2024-01-30 13:14:58 -08:00
dependabot[bot]
82dcf5a0a8 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>

[ROCm/composable_kernel commit: 84832fc42d]
 2024-01-29 09:02:52 -08:00
Illia Silin
8d92b8d30d fix CK path for hipTensor (#1143) 
[ROCm/composable_kernel commit: 4a8297c28a]
 2024-01-25 17:05:43 -08:00
rocking
133d8e762f 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

[ROCm/composable_kernel commit: 28f68a5a99]
 2024-01-25 19:53:15 +08:00
Illia Silin
b5a3426d89 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

[ROCm/composable_kernel commit: 180e572076]
 2024-01-24 13:47:48 -08:00
Bartłomiej Kocot
526905641c Fix possible linting errors in changelog (#1141) 
* Fix possible linting errors in changelog

* Update CHANGELOG.md

* Update CHANGELOG.md

* Update CHANGELOG.md

[ROCm/composable_kernel commit: 6169fbbdb3]
 2024-01-24 17:19:02 +01:00
zjing14
db32462b2d fixed return (#1138) 
[ROCm/composable_kernel commit: 1be4706366]
 2024-01-22 08:42:26 -08:00
Haocong WANG
ec7e5b1331 [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

[ROCm/composable_kernel commit: bb63b9732c]
 2024-01-19 07:02:22 -06:00
Bartłomiej Kocot
a69d7546b1 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>

[ROCm/composable_kernel commit: 7e4eb4b800]
 2024-01-19 11:29:00 +01:00
Illia Silin
6d12ce4b13 add Adam to code owners (#1136) 
[ROCm/composable_kernel commit: 38882d8ab5]
 2024-01-18 19:20:40 -06:00
randyh62
ab91e6e73e 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>

[ROCm/composable_kernel commit: 402a930a4a]
 2024-01-16 09:00:37 -08:00
Illia Silin
83d46ca80b add code owners (#1132) 
[ROCm/composable_kernel commit: c1b5b58192]
 2024-01-16 07:55:18 -08:00
Illia Silin
4612139045 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

[ROCm/composable_kernel commit: e6d099c830]
 2024-01-15 09:11:45 -08:00
dependabot[bot]
f0da02cf0c 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>

[ROCm/composable_kernel commit: 636a31015a]
 2024-01-15 09:09:13 -07:00
dependabot[bot]
483dfa6f31 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>

[ROCm/composable_kernel commit: 0ce417269d]
 2024-01-11 11:27:03 -07:00
Illia Silin
264de3901a 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

[ROCm/composable_kernel commit: 886d9eeb99]
 2024-01-09 09:43:08 -08:00
raramakr
70e816fd12 SWDEV-439954 - Use hard coded filename rather than using the macro __FILE__ for debug prints. (#1123) 
* SWDEV-439954 - Use hard coded filename rather than using the macro __FILE__ for debug prints.

Hiptensor library is using the header files from CK. Hard coded ROCm path was getting embedded into the hiptensor library, since the header file was having the macro __FILE__. Replace the macro with filename.

* fix syntax

---------

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

[ROCm/composable_kernel commit: e699dbd8a3]
 2024-01-09 08:21:47 -08:00
Illia Silin
0a6a3c7490 fix dockerfile syntax for test compilers (#1120) 
[ROCm/composable_kernel commit: 22db1e0865]
 2024-01-05 13:54:40 -08:00
randyh62
7507939ccc doc reorg and edits (#1112) 
* doc reorg and edits

* Update wrapper.rst with changes from PR #1098

* Update docs/dockerhub.rst

Co-authored-by: Bartlomiej Wroblewski <bwroblewski10@gmail.com>

* Update docs/index.rst

Co-authored-by: Bartlomiej Wroblewski <bwroblewski10@gmail.com>

* Update docs/what-is-ck.rst

Co-authored-by: Bartlomiej Wroblewski <bwroblewski10@gmail.com>

* Update docs/what-is-ck.rst

Restored to 4 bullets, with additional text for wrapper.

Co-authored-by: Bartlomiej Wroblewski <bwroblewski10@gmail.com>

* Update docs/Contributors_Guide.rst

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

* Update API_Reference_Guide.rst

using sentence case for title

* updated index structure per Lisa

* separate docker hub and tutorial

---------

Co-authored-by: Bartlomiej Wroblewski <bwroblewski10@gmail.com>
Co-authored-by: Lisa <lisajdelaney@gmail.com>
Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>

[ROCm/composable_kernel commit: a39163814e]
 2024-01-05 11:04:01 -08:00
Bartlomiej Wroblewski
421311e1fe Update the recommended version of ROCm in docs (#1110) 
[ROCm/composable_kernel commit: 61545bda35]
 2024-01-05 09:36:02 -08:00
Illia Silin
d33e0c1afc Add a docker for testing CK with rocm6.0.1 RC1. (#1119) 
* add docker for rocm6.0.1 rc1

* modify the path to clang for test compilers in CI

* fix the hipcc/clang path for test compilers in CI

* fix the dockerfile for older rocm versions

[ROCm/composable_kernel commit: d89700201b]
 2024-01-05 08:01:33 -08:00
Bartłomiej Kocot
ac9d8439ea Add missing copyrights in elementwise_permute examples (#1118) 
[ROCm/composable_kernel commit: 11e2752261]
 2024-01-04 17:38:24 +01:00
arai713
18a61d95d1 Transpose profiler fix (#1114) 
* added working example for 5D input using 1D kernel

* example with 5D input tensor and 2d kernel - not working: issues with arguments

* added updated version of 3d device op - changed descriptors/dims

* added example file to check kernel

* fixed descriptor and isSupportedArgument stride problem

* added and modified kernel for 3d - updated tids/loop

* adding some more 5d example files

* fixed some issues

* changes made for testing

* working version: fixed error in stride for A, still a bit inefficient

* cleaned up formatting/comments

* updating formatting

* more formatting fixes

* fixing cmake, adding back gpu targets in cmake script

* adding client example

* added instances for client example

* fixed errors in client example

* implemented client ex with device_elementwise.hpp and device_elementwise_3d_impl.hpp

* removed extra files

* minor formatting and naming fixes

* adding test files and profiler

* fixing minor error

* minor fix

* removed unneccesary comments, renamed files

* updated instance list for client example, added different layout example

* removing instances

* fixed error in instance generation

* remove comments

* update profiler and client example tensor layouts

* fixed errors in test/profiler

* updated vector dim access to enable vector load

* updated test/profiler files

* updated example with 1d kernel

* updating profiler

* renamed files

* disabled device op for MI300

* skip  elementwise_permute_2d on gfx94x

* Update CMakeLists.txt

* fixing CMake - disabling some GPU targets

* added transpose profiler to CMake

* fixed transpose profiler errors

* fixed instances for tests/profiler

* cleaned up code in transpose profiler source code

* added some comments, updated copyright

* made function arguments const where possible

---------

Co-authored-by: Jing Zhang <jizha@amd.com>
Co-authored-by: Jing Zhang <jizhan@amd.com>
Co-authored-by: zjing14 <zhangjing14@gmail.com>

[ROCm/composable_kernel commit: aa3e2d7967]
 2024-01-04 10:33:19 -06:00
Illia Silin
82b010e4c6 fix the cmake option syntax (#1117) 
[ROCm/composable_kernel commit: fbf31a2ea3]
 2024-01-03 07:56:44 -08:00
Bartłomiej Kocot
e69eb06fe2 Add tensor partition and generic copy for ck wrapper (#1108) 
* Add tensor partition and generic copy for ck wrapper

* Update changelog

* Stylistic fixes

* Change shape/strides logic to descriptor transforms

* Fixes

* Fix client example

* Fix comments

[ROCm/composable_kernel commit: 4234b3a691]
 2024-01-03 01:10:57 +01:00
Illia Silin
3d01c4cad6 adding -Wno-switch-default compiler flag (#1115) 
[ROCm/composable_kernel commit: b268f273de]
 2024-01-02 14:01:12 -08:00
Illia Silin
18c8f366cb change the googletest cmake syntax for older cmake versions (#1116) 
[ROCm/composable_kernel commit: 0e07dfdeab]
 2024-01-02 14:00:53 -08:00
Bartłomiej Kocot
d68c1fad45 Revert "[SWDEV-435347] disable instances failed with mainlien compiler (#1077)" (#1101) 
This reverts commit a54a4f4281.

[ROCm/composable_kernel commit: a35e466c86]
 2024-01-02 11:36:45 +01:00
Bartłomiej Kocot
35714a49ad Fix results verify in test_tensor (#1109) 
[ROCm/composable_kernel commit: 20b1ae7ced]
 2023-12-23 22:12:49 +01:00
dependabot[bot]
64713eb2d0 Bump rocm-docs-core from 0.30.2 to 0.30.3 in /docs/sphinx (#1107) 
Bumps [rocm-docs-core](https://github.com/RadeonOpenCompute/rocm-docs-core) from 0.30.2 to 0.30.3.
- [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.2...v0.30.3)

---
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>

[ROCm/composable_kernel commit: 78eb3f0b46]
 2023-12-20 14:35:25 -08:00
Artur Wojcik
e9ec2910a0 enable compilation of INSTANCES_ONLY for Windows (#1082) 
* enable compilation of INSTANCES_ONLY for Windows

* suppress ROCMChecks warnings on GoogleTests

* suppress -Wfloat-equal warning on GoogleTests

---------

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

[ROCm/composable_kernel commit: fb5bd51b42]
 2023-12-20 14:34:53 -08:00