Chao Liu
a15f7c025f
fix clang warning suppression
...
[ROCm/composable_kernel commit: ddd49ec9e7
2021-08-10 06:20:24 +00:00
Chao Liu
72c6d11a3b
tidy
...
[ROCm/composable_kernel commit: 76f3131939
2021-08-09 18:49:59 -05:00
Chao Liu
7b306ba01c
tidy
...
[ROCm/composable_kernel commit: d18428901e
2021-08-09 18:20:02 -05:00
Chao Liu
590dde14c8
tidy
...
[ROCm/composable_kernel commit: 56fc0842b3
2021-08-09 19:27:49 +00:00
Chao Liu
193f7cbbcf
tidy
...
[ROCm/composable_kernel commit: e62bae7a4a
2021-08-09 15:11:35 +00:00
Chao Liu
71a28fc501
add tidy
...
[ROCm/composable_kernel commit: 24c8728942
2021-08-08 17:41:54 +00:00
Chao Liu
b0ab2054a9
remove online compilation from CK
...
[ROCm/composable_kernel commit: ae98b52ad8
2021-08-07 00:51:05 +00:00
Chao Liu
f94e566273
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
cbd5de28e1
Deprecate static kernel ( #42 )
...
* deprecate static kernels
[ROCm/composable_kernel commit: 81c942cd7e
2021-07-08 10:40:00 -05:00
zjing14
67dcc552b6
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
817b2a47c6
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
b4dbf677ce
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
0eb214d1cd
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
3799741fee
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
08f130fed1
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
