Commit Graph

55 Commits

Author SHA1 Message Date
Chao Liu
868068eee1 implicit gemm v1r3 nchw_cyxk_nkhw
[ROCm/composable_kernel commit: a903146427]
2019-04-25 15:14:39 -05:00
Chao Liu
21988c32b4 added implicit gemm v1r3 lds_double_buffer NCHW * CYXK = KNHW, reworked static functionals
[ROCm/composable_kernel commit: 569ad66e2a]
2019-04-23 17:51:14 -05:00
Chao Liu
fe17969c81 added lds double buffer (on C dimension) for implicit gemm v1r3, as a result, it should achieve 90% of peak for all filter sizes, on CHWN format
[ROCm/composable_kernel commit: 87d8740bf5]
2019-04-19 17:49:25 -05:00
Chao Liu
26fab71682 added implicit gemm v1r3, refactored decomposition of wei tensor (loop over y, x first, and C second) to allow easy lds double buffer on C
[ROCm/composable_kernel commit: 6d066ede00]
2019-04-19 16:46:29 -05:00
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
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
8562528743 add 3x3 28x28 case
[ROCm/composable_kernel commit: 71434918bf]
2019-04-10 11:33:57 -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
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
c79fbba997 refactor
[ROCm/composable_kernel commit: b57d60c0b7]
2019-04-06 19:01:59 -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
fab54d9d58 debugging
[ROCm/composable_kernel commit: 0983d205ad]
2019-04-05 18:54:26 -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
Jing Zhang
7bea559923 add
[ROCm/composable_kernel commit: 6a3f3f951d]
2019-04-03 17:08:14 -05:00
Chao Liu
fdb7d41661 cleaning up dead code
[ROCm/composable_kernel commit: bdbc0eaad1]
2019-04-02 17:58:44 -05:00
Jing Zhang
6f8247cf59 4x4
[ROCm/composable_kernel commit: 114fdb58af]
2019-04-01 17:02:02 -05:00
Chao Liu
7cbd63b2d0 refactor
[ROCm/composable_kernel commit: e43d7bc63c]
2019-04-01 15:17:22 -05:00
Chao Liu
cd883e7581 experimenting
[ROCm/composable_kernel commit: 766b0a9eaf]
2019-03-24 12:09:57 -05:00
Chao Liu
6f1651f8a7 experimenting
[ROCm/composable_kernel commit: f35c64eb78]
2019-03-23 19:35:31 -05:00
Chao Liu
971442aab2 experiment with hip compiler
[ROCm/composable_kernel commit: 68ae0731f1]
2019-03-22 19:40:01 -05:00
Chao Liu
b24dc40c09 adding inline asm
[ROCm/composable_kernel commit: 52ae168b17]
2019-03-22 17:05:23 -05:00
Chao Liu
2c3bd06d25 Merge branch 'direct_fp16'
[ROCm/composable_kernel commit: fdaaaa500c]
2019-03-22 16:46:41 -05:00
Chao Liu
1f925812b2 hip build
[ROCm/composable_kernel commit: 8c923db423]
2019-03-22 14:22:58 -05:00
Chao Liu
9637b859d8 added int8x4
[ROCm/composable_kernel commit: e72eece8fc]
2019-03-21 09:59:40 -05:00
Chao Liu
80f2ab2e09 adding int8 direct that reads pre-vectorized data
[ROCm/composable_kernel commit: 050a1a6890]
2019-03-19 00:05:41 -05:00
Chao Liu
5ba0f64087 adding fp16 direct that reads pre-vectorized data
[ROCm/composable_kernel commit: 4f0fc72e91]
2019-03-18 15:03:17 -05:00
Chao Liu
da50d65ba0 refactoring block copy
[ROCm/composable_kernel commit: 03eef73c5b]
2019-03-17 15:36:38 -05:00
Chao Liu
0485704cb3 refactor
[ROCm/composable_kernel commit: fd8de38417]
2019-03-16 10:50:46 -05:00
Chao Liu
7bfe330532 update hip build
[ROCm/composable_kernel commit: 2c9b8c2432]
2019-03-12 17:20:11 -05:00
Chao Liu
5c5f48e664 update
[ROCm/composable_kernel commit: 0c88a3d891]
2019-03-09 13:51:08 -06:00
Chao Liu
85e6b5db6d Merge branch 'master' into implicit_gemm_fp16
[ROCm/composable_kernel commit: ce0182ce05]
2019-03-09 13:46:47 -06:00
Chao Liu
8c2635a357 refactor
[ROCm/composable_kernel commit: 7a97087713]
2019-03-09 12:59:47 -06:00
Chao Liu
005e6fe84d refactor
[ROCm/composable_kernel commit: 8edbc659b8]
2019-03-06 12:34:31 -06:00
Chao Liu
e3ab560c50 refactor
[ROCm/composable_kernel commit: 04c5527d07]
2019-03-04 17:09:20 -06:00
Chao Liu
77ec1e4016 clean up
[ROCm/composable_kernel commit: 5fd40ad768]
2019-03-02 17:27:37 -06:00
Chao Liu
7080421844 refactor
[ROCm/composable_kernel commit: 4543d17a71]
2019-02-19 22:07:15 -06:00
Chao Liu
23b7f3ef33 refactor
[ROCm/composable_kernel commit: b2b622e8b2]
2019-02-19 20:34:21 -06:00
Chao Liu
1c962a13ee device_implicit_gemm_convolution_1_chwn_csrk_khwn: use tensor copy (instead of pointwise) for writing output, 3x3 increased from 78% to 84%, 5x5 from 80% to 84%
[ROCm/composable_kernel commit: a65ef90308]
2019-02-19 11:47:46 -06:00