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
4,099 Commits 742 Branches 38 Tags
tmp-develop
Commit Graph

171 Commits

Author SHA1 Message Date
Chao Liu
51ab4abaf4 add tidy 
[ROCm/composable_kernel commit: 24c8728942]
 2021-08-08 17:41:54 +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
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
Chao Liu
459dc6cf2f Deprecate static kernel (#42) 
* deprecate static kernels

[ROCm/composable_kernel commit: 81c942cd7e]
 2021-07-08 10:40:00 -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
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
a3c89131fa Code clean up (#20) 
* tuning para,

* testing on v100

* add fp16

* remove deprecated tensor descriptor

* sync with miopen

* update build script

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

[ROCm/composable_kernel commit: 5c7cec1115]
 2020-06-23 20:31:27 -05:00
Chao Liu
4414e495ed backward data (#7) 
* enabled atomic add in tensor copy
* added gridwise GEMM
* added backward data conv using GEMM + atomic
* added backward data conv using GEMM, no atomic


[ROCm/composable_kernel commit: 8f5f64960e]
 2019-12-03 01:16:12 -06:00
Chao Liu
a9e6c3340c Refactor for MIOpen integration (#4) 
Refactor, so can bring multi-index transformation and padding support into MIOpen

[ROCm/composable_kernel commit: 52c3fe05be]
 2019-10-11 11:37:31 -05:00
Chao Liu
611944080a refactor 
[ROCm/composable_kernel commit: 21f7e9f103]
 2019-06-19 17:43:56 -05:00
Chao Liu
5f217ebda5 reorginzed files 
[ROCm/composable_kernel commit: 1566b31736]
 2019-06-13 15:12:12 -05:00
Chao Liu
11c6b2ab9a change build 
[ROCm/composable_kernel commit: c82b833d8e]
 2019-06-12 10:47:25 -05:00
Chao Liu
14451df1fd reorginze files 
[ROCm/composable_kernel commit: 81497a93a0]
 2019-06-11 23:49:51 -05:00
Chao Liu
180f7266cf updated build 
[ROCm/composable_kernel commit: fc60454e3a]
 2019-02-15 02:17:00 -06:00
Chao Liu
153629655f update build 
[ROCm/composable_kernel commit: a414e3fdf8]
 2019-02-15 02:06:34 -06:00
Chao Liu
e7f6b820cd hip build 
[ROCm/composable_kernel commit: 67c6f73ffe]
 2019-02-15 00:54:30 -06:00
Chao Liu
662ba16dc6 refactor build, clean up 
[ROCm/composable_kernel commit: e80fbbdd71]
 2019-02-14 15:10:16 -06:00
Chao Liu
6521ccba67 cpu direct conv 
[ROCm/composable_kernel commit: d51b81588f]
 2018-10-19 01:26:21 -05:00
Chao Liu
4ef894a0b8 initial build 
[ROCm/composable_kernel commit: 06c9f9fe17]
 2018-10-14 02:10:36 -05:00
Chao Liu
8bfafec554 start adding convolution 
[ROCm/composable_kernel commit: fc98757acd]
 2018-10-08 22:49:58 -05:00