ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-14 10:09:41 +00:00
Code Issues Packages Projects Releases Wiki Activity
419 Commits 742 Branches 38 Tags
b14b5d33715a2ed12ce8a554c13ee7c2addae280
Commit Graph

419 Commits

This Branch
This Branch
All Branches
Author SHA1 Message Date
Chao Liu
b14b5d3371 tidy 
[ROCm/composable_kernel commit: c3efeb5e20]
 2021-08-09 19:32:07 +00:00
Chao Liu
9c589af829 tidy 
[ROCm/composable_kernel commit: 56fc0842b3]
 2021-08-09 19:27:49 +00:00
Chao Liu
e8def0e778 tidy 
[ROCm/composable_kernel commit: 54fba515b3]
 2021-08-09 17:33:32 +00:00
Chao Liu
9e2c3c7765 tidy 
[ROCm/composable_kernel commit: e62bae7a4a]
 2021-08-09 15:11:35 +00:00
Chao Liu
51ab4abaf4 add tidy 
[ROCm/composable_kernel commit: 24c8728942]
 2021-08-08 17:41:54 +00:00
Chao Liu
cba13cb6b3 fix 
[ROCm/composable_kernel commit: 61487e0a00]
 2021-08-07 02:31:19 +00:00
Chao Liu
5ed1b840a0 remove online compilation from CK 
[ROCm/composable_kernel commit: ae98b52ad8]
 2021-08-07 00:51:05 +00:00
Chao Liu
5856acc10f refactor 
[ROCm/composable_kernel commit: cb95421311]
 2021-08-06 22:17:51 +00:00
Chao Liu
7221bedc99 Merge commit '437cc595c6e206dfebb118985b5171bbc1e29eab' into composable_kernel_init_integration_v3 
[ROCm/composable_kernel commit: 73ca970155]
 2021-08-06 21:32:27 +00:00
Chao Liu
1d8dbe3c57 Update develop (#5) (#6) 
* refactor

[ROCm/composable_kernel commit: 5781adf5cf]
 2021-08-06 16:13:07 -05:00
Chao Liu
8ce0728ae2 Merge pull request #4 from ROCmSoftwarePlatform/separate_online_compile 
Separate online compile

[ROCm/composable_kernel commit: 97e6d514f7]
 2021-08-06 16:09:22 -05:00
Chao Liu
f017e34483 refactor 
[ROCm/composable_kernel commit: 7b1ec41e5b]
 2021-08-06 20:50:01 +00:00
Chao Liu
9eb35eec87 refactor 
[ROCm/composable_kernel commit: 49c33aaea7]
 2021-08-06 19:59:53 +00:00
Chao Liu
041c48a06c rename 
[ROCm/composable_kernel commit: 54b3e73d17]
 2021-08-06 18:07:15 +00:00
Chao Liu
898807d60f add solver ConvIgemmFwdV6r1DlopsNchwKcyxNkhw; rename static ck source files 
[ROCm/composable_kernel commit: 3d32ae9404]
 2021-07-30 17:50:17 -05:00
Chao Liu
aa8c981198 Merge pull request #3 from ROCmSoftwarePlatform/format 
Update to clang-format-10

[ROCm/composable_kernel commit: f6edda6119]
 2021-07-30 17:16:23 -05:00
Chao Liu
e2352d83a9 update to clang-format-10 
[ROCm/composable_kernel commit: 82fae390fb]
 2021-07-30 16:37:00 -05:00
Chao Liu
2e187cf3ab Merge pull request #2 from asroy/master 
Update readme

[ROCm/composable_kernel commit: bd27ed6c38]
 2021-07-28 09:43:56 -05:00
Chao Liu
f3b7220822 Update README.md 
[ROCm/composable_kernel commit: 85a1429301]
 2021-07-28 09:41:38 -05:00
Chao Liu
6403529fbc Update README.md 
[ROCm/composable_kernel commit: 56f93c6f33]
 2021-07-28 09:40:44 -05:00
Chao Liu
989c345b36 Merge pull request #1 from ROCmSoftwarePlatform/some_fix_210727 
fix building issue

[ROCm/composable_kernel commit: 594f1dbe96]
 2021-07-27 13:19:11 -05:00
Chao Liu
f864700e84 fix building 
[ROCm/composable_kernel commit: 6a1bc5939c]
 2021-07-27 13:12:43 -05:00
Chao Liu
cb9222657c [MIOpen Downstream] Initial MIOpen integration (#52) 
* update online kernel wrapper bundle all descriptors in a tuple

* change __CONSTANT__ to CONSTANT

* rename

* adding tuning

* added IsValidCompileParameter

* reorginze

* adding tunable for fp16 and int8

* fix kernel compile warning and bug fixes

* suppress warning about cast CONSTANT (address space 4) pointer

* fix building issue

[ROCm/composable_kernel commit: f63a23acb1]
 2021-07-27 00:02:27 -05:00
Chao Liu
b6c15f3eec reorganize files to prepare for MIOpen integration (#51) 
* change olc cmake

* adding online compile to fwd-v4r5r2

* update scripts

* remane fwd-v4r5r2 to fwd-v6r1

* clean up

[ROCm/composable_kernel commit: 1264925422]
 2021-07-18 00:43:05 -05:00
zjing14
ca40ef6976 Add xdlops v4r4r4 into online compilation (#48) 
* init for v4r4 xdlops olc

* refactor wrap

* init impl of v4r4 nchw xdlops olc

* tuning

* test perf

* fixed v4r4 nhwc

* tuned v4r4 nhwc

* use gridwise_gemm_xdlops_v2r3

* swap a/b

* add pointer support into offline v2r3

* debugging v4r4r4 transform for olc

* change timer of olc

* refactor v4r4 xdlops nchw olc

* remove transform fun in v4r4 xdlops nhwc olc

Co-authored-by: Chao Liu <chao.liu2@amd.com>

[ROCm/composable_kernel commit: fbdf4332c7]
 2021-07-16 23:27:08 -05:00
Chao Liu
a3fcfdd2cd Change initialization method of tensor for iGEMM (#49) 
* change init method

[ROCm/composable_kernel commit: 0a72e4df94]
 2021-07-16 22:55:01 -05:00
Chao Liu
7624926d6e default iterator hack for blockwise copy (#47) 
[ROCm/composable_kernel commit: 58a8057011]
 2021-07-16 08:57:15 -05:00
Chao Liu
97e64ed534 fix bug: config for ThreadwiseDynamicTensorSliceTransfer_v2 (#46) 
[ROCm/composable_kernel commit: 1c1b56fe61]
 2021-07-09 10:13:20 -05:00
Chao Liu
d297f0d524 Create README.md (#45) 
* Create README.md

[ROCm/composable_kernel commit: 4682d070a6]
 2021-07-08 13:32:29 -05:00
Chao Liu
443ae21e78 Tweak (#44) 
* tweak

[ROCm/composable_kernel commit: aafb5eb187]
 2021-07-08 12:17:43 -05:00
Chao Liu
37fd951a2a Update default launch bounds (#43) 
* update default launch bounds

[ROCm/composable_kernel commit: 2f82cfb190]
 2021-07-08 11:26:57 -05:00
Chao Liu
459dc6cf2f Deprecate static kernel (#42) 
* deprecate static kernels

[ROCm/composable_kernel commit: 81c942cd7e]
 2021-07-08 10:40:00 -05:00
Chao Liu
0d7baf0e50 DL GEMM fp32/fp16/int8 (#41) 
* add threadwise copy the copy a tensor in one copy, added kpack to DL GEMM

* add kpack into fwd v4r5 nchw fp32

[ROCm/composable_kernel commit: b8b2d0a6d1]
 2021-07-04 22:50:29 -05:00
Chao Liu
892c52c2ed fix complain about divide by zero (#40) 
[ROCm/composable_kernel commit: 11ec07e9d1]
 2021-07-01 16:50:57 -05:00
zjing14
2331d228e2 xdlops_v4r4_fwd fp32/fp16 (#34) 
* create files for xdlops

* working on blockwise_gemm_xdlops

* add KReduction

* add m/n repeats

* add 2x2 pipeline

* added 128x128 wavegemm

* use StaticBuffer of vector_type

* break vector type to blk_size

* add kpack into xldops_gemm and blockwise_gemm

* abroadcast only

* add fp32 mfma instructions

* adding fp16 mfma

* pack half4_t

* rename kperwave to kpack

* add 32x32x8fp16

* add fp16 mfma

* clean code

* clean code

* V4r4 xdlops kpack (#35)

* add kpack with incorrect results

* bug fix for make_dynamic_naive_tensor_descriptor_aligned_v2

* add 1x1 kernel

* add gridwise_gemm_v2 - single_buffer

* enabled dwordx4 for fp16

Co-authored-by: Chao Liu <chao.liu2@amd.com>

* refactor fwd-v4r4-xdlops

* add v4r4-nhwc-xdlop

* improve some perf of nhwc and nchw by tuning parameters, and change scheuduling in gridwise-gemm loop

* tweak scheduling in gridwise gemm

* add v4r3 with a single output copy

* init commit: output with slice win

* adding sliceWin

* add multiple repeats pattern

* starting adding bwd-v4r1-xdlops

* use tuple as SrcBuffer

* adding bwd-data v4r1 nhwc xdlops

* fix bug in make_dynamic_naive_tensor_descriptor_aligned_v2()

* fix bug in host bwd-data conv

* initial implementation of bwd-data v4r1 nhwc xdlops

* add launch bound flags

* enable launch bound

* add m/nrepeat=4

* tweak bwd-data v4r1 nhwc xdlops

* added bwd-data v4r1 nhwc xlops with output A and weight B

* add fwd-v4r4 nhwc xdlops, A input, B weight, C output

Co-authored-by: Chao Liu <chao.liu2@amd.com>

[ROCm/composable_kernel commit: 3835318cc3]
 2021-07-01 14:33:00 -05:00
Qianfeng
0d278b8cc8 Add online compilation for dynamic kernels (#37) 
* Add online-compiling facility

* Synchronize from fwd-v4r5 and implement host interfaces to call conv-fwd v4r4/v4r5 using on-line compiling method

* Tiny adjustment to time reporting

* Use object assignment to replace explicit bytes copying in the first kernel of v4r4/v4r5

* Use single thread to assign descriptor object to device memory

* Adjust to the workload assignment of the two kernels of v4r4 (experimental)

* Revert "Adjust to the workload assignment of the two kernels of v4r4 (experimental)"

This reverts commit eb38461456bb0c82b6c0d32cdd616e181907e20c.

* Update to make constexpr for generating descriptor types in kernel 2 of dynamic conv-fwd v4r4

* Update to dynamic conv-fwd v4r4 online-compiling

* Update to dynamic conv-fwd v4r5 online-compiling (result not accurate)

* Tiny update to driver/CMakeLists.txt

* clang-format

* Tiny comments change

* Add env OLC_DUMP_SAVE_TMP_DIR to support saving of temperary dir

* Fwd v4r5 olc perf (#39)

* added hip-clang flags that fix perf issue of online compilation

* fix bug for olc fwd-v4r5-nchw

* Move constexpr and type reference statements out of the function body in conv-fwd v4r4/v4r5 kernel wrapper

* Remove printing in hip_build_utils.cpp

* Update to root CMakeLists.txt

* Revert "Move constexpr and type reference statements out of the function body in conv-fwd v4r4/v4r5 kernel wrapper"

This reverts commit 3d2c5d8ecdd8298b72d127110500ed5b38d9835c.

Co-authored-by: Chao Liu <chao.liu2@amd.com>
Co-authored-by: Chao Liu <lc.roy86@gmail.com>
Co-authored-by: root <root@dc-smc-18.amd.com>

[ROCm/composable_kernel commit: 1685048a67]
 2021-06-24 08:34:19 -05:00
Chao Liu
56fef2f355 pass-by-void-pointer for gridwise_dynamic_gemm_v1r2 (#38) 
* pass-by-void-pointer for gridwise_dynamic_gemm_v1r2

* use pass-by-value by default

[ROCm/composable_kernel commit: d2315b0dfc]
 2021-06-19 13:43:45 -05:00
Chao Liu
c55129e8f5 Restructure gridwise and blockwise GEMM, add tensor contraction and FWD-v4r5 (#36) 
* experimenting magic number division

* overhauling fwd-v4r4 to clearly reflect transformation graph

* added fwd-v4r5

* bug fix for make_dynamic_naive_tensor_descriptor_aligned_v2

* bug fix and added sanity-check in transform_dynamic_tensor_descriptor

* added conv_driver_v2

[ROCm/composable_kernel commit: 30072aec37]
 2021-06-09 23:53:08 -05:00
Chao Liu
d87338c208 reorganize some files (#33) 
[ROCm/composable_kernel commit: 71d6b19d18]
 2021-05-12 14:15:38 -05:00
Chao Liu
0ac401f4f3 Use DynamicBuffer instead of raw pointer (#32) 
* Use DynamicBuffer to hold raw pointer (to global and LDS memory)

* add workaround for compiler issue (inefficient ISA) of ds_write for int8x4, int8x8, int8x16

[ROCm/composable_kernel commit: 78b987fbd6]
 2021-05-12 13:10:42 -05:00
Chao Liu
e100ee5732 No raw index calculation (#31) 
* Replace most raw index calculation to coordinate transformation
* Overhaul blockwise and threadwise GEMM
* Overhaul driver for gridwies GEMM kernel

Co-authored-by: Jing Zhang <jizhan@amd.com>

[ROCm/composable_kernel commit: 01055d95d9]
 2021-05-11 00:09:25 -05:00
Chao Liu
c67332b930 Use Tuple and vector_type instead of Array for holding tensor data (#30) 
* replacing array with tuple and vector for tensor data

[ROCm/composable_kernel commit: d075adf126]
 2021-04-28 13:10:33 -05:00
Chao Liu
2501a44530 Overhaul vector_type and use real vector for int8x4_t instead of aliasing from int32_t (#29) 
* overhaul vector_type, make int8x4_t real vector instead of aliasing from int32_t

[ROCm/composable_kernel commit: e4790c250c]
 2021-04-12 23:48:43 -05:00
Chao Liu
4626ccca4e Initial implementation of magic number division and "Merge" transformation that use it (#28) 
* initial implementation for magic number division and DynamicMerge_v2_magic_division that uses it

* turn off DynamicMerge_v2_magic_division that use magic number division by default

[ROCm/composable_kernel commit: 3bf52e60c5]
 2021-04-12 21:32:55 -05:00
zjing14
2457224dc9 Hybrid direct + implicit GEMM forward convolution NCHWc v5r1 (#25) 
* Hybrid direct + implicit GEMM forward convolution NCHWc v5r1. Input tensor bypass LDS. Support fp32/fp16/int8

[ROCm/composable_kernel commit: 792a20fa5b]
 2021-04-07 16:47:29 -05:00
Chao Liu
ca8a932775 Fix performance issue when passing tensor descriptor from host to kernel by void pointers (#27) 
* use address_space(4) in kernel signature to fix performance issue when passing tensor descriptor from host to kernel by (void) pointers

* remove passing by pointer* option (only use pass by value or void*)

[ROCm/composable_kernel commit: d2217f3040]
 2021-04-06 17:49:57 -05:00
zjing14
404ce5308a bug fix for buffer resource setting (#26) 
[ROCm/composable_kernel commit: 6a5ea49309]
 2021-04-06 16:59:52 -05:00
Chao Liu
e2753e68bd Dynamic tensor descriptor (#24) 
* support dynamic tensor descriptor

* use buffer load OOB feature for padding case

* add navi support

* add int8x4 inference kernel

Co-authored-by: Chao Liu <chao@ixt-rack-81.local.lan>
Co-authored-by: Jing Zhang <jizhan@amd.com>

[ROCm/composable_kernel commit: fcbb978828]
 2021-03-25 13:51:11 -05:00
Chao Liu
821c1ff9b9 Bwd Data NHWC (#22) 
* fix buffer_store bug
* remove obsolete kernels
* add bwd-data-v5r1-nhwc 

[ROCm/composable_kernel commit: bbcb67d0aa]
 2020-08-06 12:22:11 -05:00
Chao Liu
ceffe74a0a Improve buffer address for out of bound check (#21) 
* Use buffer load built-in OOB check. buffer size is limited to 2GB.
* buffer APIs use combined wave and thread offset
* use uint32_t for addr shift in buffer addressing

[ROCm/composable_kernel commit: ac62d13ecd]
 2020-07-29 18:04:09 -05:00