Commit Graph

190 Commits

Author SHA1 Message Date
Chao Liu
f367814ef7 added GridwiseConvolutionImplicitGemm_v1r2_nchw_cyxk_khwn
[ROCm/composable_kernel commit: 5ce19234a4]
2019-04-19 14:22:02 -05:00
Chao Liu
d0244d3a51 implicit gemm v1r2: adding support for nchw
[ROCm/composable_kernel commit: 19f17df47a]
2019-04-18 11:49:09 -05:00
Chao Liu
482e5e9293 refactor ConstantTensorDescriptor and functional
[ROCm/composable_kernel commit: 17f3d2d4bc]
2019-04-16 17:36:18 -05:00
Chao Liu
fb7f132a64 refactor
[ROCm/composable_kernel commit: a2cf803c7e]
2019-04-13 16:21:55 -05:00
Chao Liu
70e58e6331 refactor
[ROCm/composable_kernel commit: 3ffb824ed3]
2019-04-13 16:21:44 -05:00
Chao Liu
59e72ade87 tuning
[ROCm/composable_kernel commit: 7d8daba741]
2019-04-13 15:53:03 -05:00
Chao Liu
8b7eafe959 implicit gemm v1r2: only load 1d filter
[ROCm/composable_kernel commit: 00899f191b]
2019-04-13 11:19:17 -05:00
Chao Liu
2fc34a9169 tuned implicit gemm v1 for 3x3 on AMD to 82%. Fixed a bug in 4d tensor blockwise copy.
[ROCm/composable_kernel commit: 96ee9571e2]
2019-04-10 18:10:18 -05:00
Chao Liu
c8aed9eef4 update flops calculation
[ROCm/composable_kernel commit: edc89778c3]
2019-04-10 15:35:00 -05:00
Chao Liu
3d62f618d4 simplify blockwise batched GEMM
[ROCm/composable_kernel commit: 5696c81ffd]
2019-04-10 15:29:35 -05:00
Chao Liu
8562528743 add 3x3 28x28 case
[ROCm/composable_kernel commit: 71434918bf]
2019-04-10 11:33:57 -05:00
Chao Liu
922ace65cd clean up
[ROCm/composable_kernel commit: d86a5e4b47]
2019-04-09 19:12:16 -05:00
Chao Liu
35140bbc9e enabled ds_read_b128 and ds_write_b128 on hip c++
[ROCm/composable_kernel commit: e624df922d]
2019-04-09 19:05:44 -05:00
Chao Liu
75ca00f748 tidy yp
[ROCm/composable_kernel commit: 471830a052]
2019-04-09 18:07:36 -05:00
Chao Liu
4fba215ede refactor
[ROCm/composable_kernel commit: 1bd880a674]
2019-04-09 16:13:14 -05:00
Chao Liu
e99510c4db load smaller weight tensor
[ROCm/composable_kernel commit: 796f72e26e]
2019-04-08 23:57:13 -05:00
Chao Liu
85a7489a14 refactor
[ROCm/composable_kernel commit: 5b36aeadeb]
2019-04-08 16:13:17 -05:00
Chao Liu
709d9581bf added implicit_gemm_v1 lds double_buffer
[ROCm/composable_kernel commit: cc0fa73acd]
2019-04-08 14:11:55 -05:00
Chao Liu
217b1306e6 add more assertion
[ROCm/composable_kernel commit: c075d3f7d9]
2019-04-08 12:02:56 -05:00
Chao Liu
8cb0a5b05f tidy up
[ROCm/composable_kernel commit: 268d1c717c]
2019-04-08 10:48:29 -05:00
Chao Liu
d430879858 debugging implicit gemm v1: use 10d tensor output
[ROCm/composable_kernel commit: c9fa46af0b]
2019-04-08 10:27:32 -05:00
Chao Liu
14d997382a refactor
[ROCm/composable_kernel commit: 90abf42799]
2019-04-06 19:39:58 -05:00
Chao Liu
c79fbba997 refactor
[ROCm/composable_kernel commit: b57d60c0b7]
2019-04-06 19:01:59 -05:00
Chao Liu
302100002b use dedicated threadwise_copy for 1x1, perf at 80%
[ROCm/composable_kernel commit: bd0098afb3]
2019-04-06 18:40:54 -05:00
Chao Liu
20df62cc31 clean up
[ROCm/composable_kernel commit: 5245a0162b]
2019-04-06 16:27:07 -05:00
Chao Liu
2cd5fbe227 debugged: CUDA should use its own float4 definition
[ROCm/composable_kernel commit: 7a251a0922]
2019-04-06 15:44:53 -05:00
Chao Liu
a3f850c5e6 debugging
[ROCm/composable_kernel commit: f6cb5b846d]
2019-04-06 15:10:40 -05:00
Chao Liu
fab54d9d58 debugging
[ROCm/composable_kernel commit: 0983d205ad]
2019-04-05 18:54:26 -05:00
Chao Liu
173ae61433 preload don't wait
[ROCm/composable_kernel commit: bae2333791]
2019-04-05 02:42:54 -05:00
Chao Liu
6d76a2d59a clipboard float4 copy and paste C++ code
[ROCm/composable_kernel commit: dabfa77fc6]
2019-04-05 02:13:29 -05:00
Chao Liu
008c584abb Merge branch 'master' into inline_asm_v2
[ROCm/composable_kernel commit: 605afd0fb6]
2019-04-04 18:40:23 -05:00
Chao Liu
ed1c49376e Merge branch 'inline_asm_v2' of github.com:asroy/modular_convolution into inline_asm_v2
[ROCm/composable_kernel commit: 66edb2590d]
2019-04-04 17:37:13 -05:00
Chao Liu
1b62e612d9 unroll even-odd loop
[ROCm/composable_kernel commit: 19b4179798]
2019-04-04 17:36:02 -05:00
Jing Zhang
9c54da52e8 clean code
[ROCm/composable_kernel commit: 62c4d5dff3]
2019-04-04 14:52:42 -05:00
Jing Zhang
43283f6452 unroll k
[ROCm/composable_kernel commit: 313f3c07d2]
2019-04-04 11:43:37 -05:00
Jing Zhang
94179601ed add debug
[ROCm/composable_kernel commit: 0f620a9018]
2019-04-04 11:20:43 -05:00
Chao Liu
4b71bff0ec add asm into lds_double_buffer version
[ROCm/composable_kernel commit: fbc7817bbb]
2019-04-04 10:38:49 -05:00
Jing Zhang
5fce583246 optimize global_load
[ROCm/composable_kernel commit: 155d78594b]
2019-04-03 21:47:58 -05:00
Chao Liu
6bf787f524 enable 128x128 block gemm
[ROCm/composable_kernel commit: 05d7a0875c]
2019-04-03 19:04:17 -05:00
Jing Zhang
7bea559923 add
[ROCm/composable_kernel commit: 6a3f3f951d]
2019-04-03 17:08:14 -05:00
Jing Zhang
1c2ae33d1d fixed wait
[ROCm/composable_kernel commit: b188c0d243]
2019-04-03 15:30:19 -05:00
Jing Zhang
d0996d2840 increase depth of pipeline
[ROCm/composable_kernel commit: 0de4286a4f]
2019-04-03 15:19:40 -05:00
Jing Zhang
efe1f1bff1 add inline asm
[ROCm/composable_kernel commit: 6fef303ec6]
2019-04-03 14:36:19 -05:00
Jing Zhang
73661b97c0 inline asm
[ROCm/composable_kernel commit: 0d6aa311e9]
2019-04-03 14:34:41 -05:00
Jing Zhang
1c8570bdf5 refactor inline asm
[ROCm/composable_kernel commit: 753b98b59a]
2019-04-03 13:23:56 -05:00
Chao Liu
770c141f35 add script to extrac asm on hip
[ROCm/composable_kernel commit: 6166233e05]
2019-04-03 10:36:18 -05:00
Chao Liu
e277457dce tidy up
[ROCm/composable_kernel commit: e2313c9eca]
2019-04-02 20:30:00 -05:00
Chao Liu
9d61f2597a add cuda extract_asm script
[ROCm/composable_kernel commit: e6c86f81b5]
2019-04-02 20:26:58 -05:00
Chao Liu
0b2a76553d add a missing file
[ROCm/composable_kernel commit: c234741045]
2019-04-02 20:19:24 -05:00
Chao Liu
e423954e6e puting gridwise convolution into its own class
[ROCm/composable_kernel commit: 6290e0b080]
2019-04-02 20:18:01 -05:00