Commit Graph

170 Commits

Author SHA1 Message Date
Chao Liu
2a43644437 Merge remote-tracking branch 'origin/develop' into miopen_downstream-dynamic_reduction_pr
[ROCm/composable_kernel commit: b725e3fc84]
2021-09-21 11:55:26 -05:00
Chao Liu
f1d7806427 Add a version of Merge transform that use integerdivision and mod (#25)
* add Merg_v3_division_mod

* refactor

[ROCm/composable_kernel commit: f3acd2510b]
2021-09-05 12:57:57 -05:00
ltqin
2f4f6427f5 Backward weight v4r4r2 with xdlops (#18)
* start

* modify transformat

* modify device convolutiion

* modify host

* added host conv bwd and wrw

* remove bwd, seperate wrw

* clean

* hacall k to zero

* out log

* fixed

* fixed

* change to (out in wei)

* input hack

* hack to out

* format

* fix by comments

* change wei hacks(wei transform has not merge)

* fix program once issue

* fix review comment

* fix vector load issue

* tweak

Co-authored-by: ltqin <letaoqin@amd.com>
Co-authored-by: Jing Zhang <jizhan@amd.com>
Co-authored-by: Chao Liu <chao.liu2@amd.com>

[ROCm/composable_kernel commit: 627d8ef35a]
2021-08-30 22:49:17 -05:00
Chao Liu
a44dd0d851 Misc fixes (#24)
* use cast_pointer_to_generic_address_space() in v6r1 kernel wrapper, DynamcBuffer and buffer_load take customized invalid-element-value, add buffer_load/store for fp64

* use remove_cvref_t

[ROCm/composable_kernel commit: 10bb811060]
2021-08-26 20:05:19 -05:00
Qianfeng
b315c39b11 [SWDEV-281541][MSRCHA-100] Implementation of Dynamic Generic Reduction (#1108)
* add solver ConvIgemmFwdV6r1DlopsNchwKcyxNkhw; rename static ck source files

* make inner product compatible on gfx900

* Update src/include/miopen/solver/ck_utility_common.hpp

* compiler parameter use stream

* use int instead of index_t in kernel wrapper

* DynamicBuffer, StaticBuffer, amd_buffer_load support customized value for invalid element

* Add dynamic generic reduction kernel layer (kernel wrappers, kernel implementations and utilities)

* Some updates to dynamic composable kernel facility for the need of dynamic generic reduction

* Update to generic reduction C++ host interface layer to support dynamic generic reduction

* Update to remove tidy complaints in host interface layer

* Change the unary operator form from void op(T &x) to T op(T x)

* Update to pass single workspace pointer for all kernels (fix for OpenCL backend)

* Use cppcheck-suppress to prevent some strange warnings

* Re-use operator [] and () for DynamicBuffer and update to depending codes

* Remove useless codes in first call threadwise/warpwise/blockwise kernel wrappers

* [performance] Remove un-needed local buffer initialization

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

[ROCm/composable_kernel commit: 9e80cdceb7]
2021-08-26 18:04:55 -07:00
zjing14
20e76381da GlobalAtomicAdd for fp32/int32 (#23)
* add f32/i32 atomicAdd support into dynamicBuffer, and enable it in v1r3

* fixed

* fixed

* update comment

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

[ROCm/composable_kernel commit: a7a758d8ce]
2021-08-25 10:55:55 -05:00
zjing14
14e4d7ba7d Xdlops refactor fix (#22)
* added constexpr ahead of adptor; clean unused driver; rename M/NPerWave to M/NPerXDL

* fixed bwd

* fixed comment

[ROCm/composable_kernel commit: 9d3f634a3c]
2021-08-23 11:22:10 -05:00
Chao Liu
eda5a8852c magic division use __umulhi() (#19)
[ROCm/composable_kernel commit: c6f26bb480]
2021-08-23 10:40:27 -05:00
Chao Liu
ee428d2d6f Composable kernel init integration v3 (#1097)
* Squashed 'src/composable_kernel/' content from commit f6edda611

git-subtree-dir: src/composable_kernel
git-subtree-split: f6edda6119

* add solver ConvIgemmFwdV6r1DlopsNchwKcyxNkhw; rename static ck source files

* Squashed 'src/composable_kernel/' changes from f6edda611..5781adf5c

5781adf5c Update develop (#5) (#6)
97e6d514f Merge pull request #4 from ROCmSoftwarePlatform/separate_online_compile
7b1ec41e5 refactor
49c33aaea refactor
54b3e73d1 rename

git-subtree-dir: src/composable_kernel
git-subtree-split: 5781adf5cf

* fix

* refactor

* remove online compilation from CK

* refactor

* fix

* add ctest

* add c-style pointer cast

* vector/scalar pointer cast use c-style pointer cast instead of reinterpret_cast

* fix clang warning suppression

* tidy

* suppress cppcheck

* fix enum issue

* revert chagnes to hip build

* fix kernel filename

* update CK build script

* rename

* rename

* make innner product compatiable on gfx900

* Update src/include/miopen/solver/ck_utility_common.hpp

Co-authored-by: JD <Jehandad.Khan@amd.com>

* compiler parameter use stream

* use int instead of index_t in kernel wrapper

* DynamicBuffer, StaticBuffer, amd_buffer_load support customized value for invalid element

* refactor

* refactor

* change cmakelist

* change ck common utility

* fix

Co-authored-by: JD <Jehandad.Khan@amd.com>

[ROCm/composable_kernel commit: 6fe3627a9e]
2021-08-19 10:55:03 -05:00
zjing14
6dde0995c5 refactor dynamic xdlops iGemm (#13)
* xdlops refactor

* fixed commnt

* clean xdlops_gemm

* add make c into xldops-gemm

* change mfma_info

* refactor xdlops, hide c desc

* clean

* clean

* clean

* apply hacks changes to v4r4r4_nhwc

* rename hacks and use single stage adapter

* enable fp16 mfma

[ROCm/composable_kernel commit: a2ad6d3531]
2021-08-19 09:54:10 -05:00
Chao Liu
e25c4c2f15 refactor
[ROCm/composable_kernel commit: 67ad47e7c1]
2021-08-16 21:01:33 +00:00
Chao Liu
27048b7714 refactor
[ROCm/composable_kernel commit: 16effa767c]
2021-08-16 20:36:47 +00:00
Chao Liu
65e8349051 DynamicBuffer, StaticBuffer, amd_buffer_load support customized value for invalid element
[ROCm/composable_kernel commit: a91b68dfcb]
2021-08-13 23:40:19 +00:00
Chao Liu
b3759bf6a2 use int instead of index_t in kernel wrapper
[ROCm/composable_kernel commit: 2cbabbba54]
2021-08-13 20:55:39 +00:00
Chao Liu
9f40048d1c make innner product compatiable on gfx900
[ROCm/composable_kernel commit: f2ac7832c6]
2021-08-11 09:42:53 -05:00
Chao Liu
f7df8c7eea rename
[ROCm/composable_kernel commit: 4e57b30a6a]
2021-08-11 00:08:42 +00:00
Chao Liu
1e312fef12 rename
[ROCm/composable_kernel commit: c03045ce2d]
2021-08-10 23:45:36 +00:00
Chao Liu
594b1cf918 fix enum issue
[ROCm/composable_kernel commit: d626dccc95]
2021-08-10 20:55:13 +00:00
Chao Liu
a7c943abab fix clang warning suppression
[ROCm/composable_kernel commit: ddd49ec9e7]
2021-08-10 06:20:24 +00:00
Chao Liu
d49e0ddcb2 vector/scalar pointer cast use c-style pointer cast instead of reinterpret_cast
[ROCm/composable_kernel commit: 4f566c6221]
2021-08-10 05:55:20 +00:00
Chao Liu
314b9d78ef add c-style pointer cast
[ROCm/composable_kernel commit: 172036d728]
2021-08-10 00:01:52 -05:00
Chao Liu
d4b35bd09f tidy
[ROCm/composable_kernel commit: 76f3131939]
2021-08-09 18:49:59 -05:00
Chao Liu
cb2edf2100 tidy
[ROCm/composable_kernel commit: d18428901e]
2021-08-09 18:20:02 -05:00
Chao Liu
4771cfa340 tidy
[ROCm/composable_kernel commit: f885c131d8]
2021-08-09 22:13:47 +00:00
Chao Liu
eb7f9f35b2 tidy
[ROCm/composable_kernel commit: 80120f0a0c]
2021-08-09 21:10:09 +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
51ab4abaf4 add tidy
[ROCm/composable_kernel commit: 24c8728942]
2021-08-08 17:41:54 +00:00
Chao Liu
e2352d83a9 update to clang-format-10
[ROCm/composable_kernel commit: 82fae390fb]
2021-07-30 16:37:00 -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
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
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