Skip to content

[conv bwd-weight]Binding gemm k1 to conv n#202

Merged
asroy merged 48 commits into
developfrom
binding_gemm_k1_to_conv_n
May 20, 2022
Merged

[conv bwd-weight]Binding gemm k1 to conv n#202
asroy merged 48 commits into
developfrom
binding_gemm_k1_to_conv_n

Conversation

@shaojiewang
Copy link
Copy Markdown
Contributor

@shaojiewang shaojiewang commented Apr 25, 2022 •
edited
Loading

conv bwd-weight optimization:

  1. reduce move slice window instructions by binding gemm k1 to conv n dimension.
  2. obvious performance improvement for big tile size.

shaojiewang and others added 30 commits April 14, 2022 10:16
@shaojiewang
add some instance to develop
f9fde06
@shaojiewang
avoid bank conflicts for wrw for all instance
09f365a
@shaojiewang
add small K1 test
0a1e41a
@shaojiewang
delete some unused instance
0fd4df3
@shaojiewang
binding gemm k1 to conv n
45440f1
@shaojiewang
try using half_4 to do ds_read
ff6aaca
@shaojiewang
reset buffer load oob and ds memcpy to default option
230a41c
@shaojiewang
remove useless instances
81ffce2
@shaojiewang
remove redandunt space
a6ebdb4
@shaojiewang
remove printf code
c58fc51
@shaojiewang
Merge branch 'develop' into wrw_conv_impr
fc17eb4
@shaojiewang
clang-format-10 change
1654fcc
@shaojiewang
use fastest config
e6b32ff
@shaojiewang
fix clang format for the other files
260dcdb
@shaojiewang
remove gemmk0 pad for output
add7421
@shaojiewang
add gemmk padding macro
798e93e
@shaojiewang
add bank length computation
93871ca
@shaojiewang
add template to distinguish the instance that need lds padding for wrw
1e5c712
@shaojiewang
use rocm5.1 as docker
eb09227
@shaojiewang
Merge branch 'develop' into wrw_conv_impr
2e6eaf6
use integer value for GEMM test
579e8e7
@shaojiewang
add Right padding macro
3e7f799
Merge remote-tracking branch 'origin/develop' into wrw_conv_impr
507e149
Merge remote-tracking branch 'origin/fix_test' into wrw_conv_impr
60ee26a
@shaojiewang
Merge branch 'develop' into wrw_conv_impr
5c8ad21
@shaojiewang
merge dev branch
19bb36d
@shaojiewang
add 2 test asm code
ebf3d70
@shaojiewang
using 256x256x32 tile size
438138c
@shaojiewang
1. move dedicated transform into gridwisegemm's head file. 2. make ld…
28bb628 
…s tensor params a struct templete. 3. remove useless code
@shaojiewang
using small vec
d4b7361
@asroy asroy added the WIP label May 7, 2022
@shaojiewang shaojiewang marked this pull request as ready for review May 13, 2022 12:31
@shaojiewang shaojiewang requested a review from asroy May 13, 2022 12:32
@shaojiewang shaojiewang changed the title [WIP][Wrw]Binding gemm k1 to conv n [conv bwd-weight]Binding gemm k1 to conv n May 13, 2022
@shaojiewang shaojiewang removed the WIP label May 13, 2022
@zjing14
Copy link
Copy Markdown
Contributor

zjing14 commented May 16, 2022

@asroy Review required

@asroy asroy merged commit 070619f into develop May 20, 2022
@junliume junliume deleted the binding_gemm_k1_to_conv_n branch October 21, 2023 06:10
hyoon1 pushed a commit to hyoon1/composable_kernel that referenced this pull request Mar 19, 2026
@tridao
Merge pull request ROCm#202 from fedebotu/main
5bf7f57 
[BugFix] avoid bug on ImportError
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

1 more reviewer

@asroy asroy asroy approved these changes

Reviewers whose approvals may not affect merge requirements

Assignees

No one assigned

Labels

CI - Pass

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

3 participants

@shaojiewang @zjing14 @asroy