Commit Graph

21 Commits

Author SHA1 Message Date
Chao Liu
58864d1082 fix bug in gridwise gemm xdlops v2r3 (#45)
[ROCm/composable_kernel commit: d5297abae9]
2021-10-21 16:42:24 -05:00
ltqin
3341ddbdf5 add nchw atomic , nhwc and nhwc atomic method for backward weight (#30)
* add add new algorithm from v4r4r2

* program once issue

* add split k functiion

* redefine code

* add a matrix unmerge

* add b matrix unmerge k0

* trans a and b to gridegemm

* nhwc init

* no hacks and vector load

* add hacks

* modify some parameter

* fix tuning prometer for fp32

* fix tuning prometer for fp16

* start change gridwise k split

* init ok

* revome a b matrix k0mk1 desc in grid

* carewrite lculate gridsize

* add kbatch to CalculateBottomIndex

* remove some unused funtion

* add clear data function before call kernel

* out hacks

* in hacks

* rename device convolution file and function name

* modify kBatch value

* fix some tuning code

* start from v4r4 nhwc

* nhwc atomic is able to run

* just for fp32

* enable nchw atomic

* tweak

* tweak

* re-arrange gridwise gemm hot loop for wrw

* add wrw v4r5

* v4r4r5 fp16

* v4r4r4 fp16

* v4r4r2 fp16

* V4R4R4XDLNHWC fp16

* V4R4R2XDLATOMICNCHW fp16

* adjust for fp16

* input gridsize

* change kbatch to gridsize

* testing wrw

* clean up

* k_batch to gridsize

* fix bug

* wrw v4r4r4 kbatch change to gride size

* wrw v4r4r2 kbatch change to gride size

* after merge , change gridwise gemm v2r4

* change MakeCBlockClusterAdaptor

* other method use new gridwise gemm

* clean up

* chapad method nge to make_right_pad_transform

* kbatch out from transform function

* clean up and fix bug

* fix bug

* using function type reduce template parameters

* using auto replace define fuction type

* clean up

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

[ROCm/composable_kernel commit: fd49ff8080]
2021-10-19 18:42:34 -05:00
Chao Liu
0e75841071 Tweak GEMM kernel (#38)
* add parameters

* tweak gemm

* tweak

* update conv

* update script

* adding bwd 1x1

* update script

* adding 1x1 bwd

* debugging bwd 1x1 failure

* update script

* update script

* test

* test v100

* clean up

[ROCm/composable_kernel commit: b3e8d57d51]
2021-10-06 11:12:36 -05:00
Chao Liu
115f77e17a GEMM driver and kernel (#29)
* add gemm driver

* tweak

* add gemm kernel: mk_kn_mn and km_kn_mn

* tweak

* add GEMM km_nk_mn

* fix comment

[ROCm/composable_kernel commit: 19613902b5]
2021-09-05 12:41:28 -05:00
ltqin
79b671c5dd 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
zjing14
6d56ee7d00 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
zjing14
3cacb0c037 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
c5ea28dc43 make innner product compatiable on gfx900
[ROCm/composable_kernel commit: f2ac7832c6]
2021-08-11 09:42:53 -05:00
Chao Liu
0bbcd05e4e rename
[ROCm/composable_kernel commit: 4e57b30a6a]
2021-08-11 00:08:42 +00:00
Chao Liu
c5a4edb9e8 rename
[ROCm/composable_kernel commit: c03045ce2d]
2021-08-10 23:45:36 +00:00
Chao Liu
a0df0eb029 fix kernel filename
[ROCm/composable_kernel commit: 2c48039d0e]
2021-08-10 22:15:23 +00:00
Chao Liu
a15f7c025f fix clang warning suppression
[ROCm/composable_kernel commit: ddd49ec9e7]
2021-08-10 06:20:24 +00:00
Chao Liu
ab465fca4c 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
17b084ce2c add c-style pointer cast
[ROCm/composable_kernel commit: 172036d728]
2021-08-10 00:01:52 -05: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
eb778cb91b tidy
[ROCm/composable_kernel commit: f885c131d8]
2021-08-09 22:13:47 +00:00
Chao Liu
590dde14c8 tidy
[ROCm/composable_kernel commit: 56fc0842b3]
2021-08-09 19:27:49 +00:00
Chao Liu
7885261dc6 tidy
[ROCm/composable_kernel commit: 54fba515b3]
2021-08-09 17:33:32 +00:00
Chao Liu
e02d6a0f21 [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
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