ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-11 17:00:18 +00:00
Code Issues Packages Projects Releases Wiki Activity
368 Commits 739 Branches 38 Tags
7d09790a0ac2be6c150d25654fdd9d05d392b34f
Commit Graph

368 Commits

This Branch
This Branch
All Branches
Author SHA1 Message Date
Chao Liu
7d09790a0a MIOpen integration (#15) 
* renaming
 2020-02-18 10:42:18 -06:00
Chao Liu
1a66e35b6f MIopen integration (#13) 
* update for miopen integration: cosmetic refactor
 2020-02-17 09:53:20 -06:00
Chao Liu
3406a1148a Update for recent MIOpen integration (#11) 
* update for MIOpen integration
 2020-01-27 15:29:33 -06:00
Chao Liu
c5da0377fb Added bwd data v3r1 v4r1, tweaking v1 (#10) 
* Added bwd data v3r1: breaking down compute into a series of load balanced GEMM, and launch in a single kernel
* Added bwd data v4r1: like v3r1, but launch GEMMs in multiple kernels
* Tweaked v1r1  and v1r2 (atomic) on AMD GPU
 2020-01-20 10:20:03 -06:00
Chao Liu
e2b4c5b469 update implicit GEMM forward v4r4 to use gridwise gemm (#9) 
* updated fwd v4r4 to use gridwise gemm
* updated gridwise gemm api calls in bwd-data v1r1 and v2r1
 2019-12-05 12:36:36 -06:00
Chao Liu
19a93dac05 fixed faulty padding API calls (#8) 2019-12-03 01:46:44 -06:00
Chao Liu
8f5f64960e backward data (#7) 
* enabled atomic add in tensor copy
* added gridwise GEMM
* added backward data conv using GEMM + atomic
* added backward data conv using GEMM, no atomic
 2019-12-03 01:16:12 -06:00
Chao Liu
31ded4ac4b remove dead file (#6) 2019-11-04 17:13:38 -06:00
Chao Liu
562e1e2767 MIOpen integration: recent bug fixes from MIOpen (#5) 2019-11-04 16:51:12 -06:00
Chao Liu
52c3fe05be Refactor for MIOpen integration (#4) 
Refactor, so can bring multi-index transformation and padding support into MIOpen
 2019-10-11 11:37:31 -05:00
Chao Liu
9aaeacc82b Merge pull request #3 from asroy/clean_up 
enable type conversion in ThreadwiseGenericTensorSliceCopy_v2r1 and BlockwiseGenericTensorSliceCopy_v2
 2019-09-30 15:14:02 -05:00
Chao Liu
cf21818455 enable type conversion in blockwise copy v2 and threadwise copy v2r1 2019-09-30 15:11:05 -05:00
Chao Liu
012d3a071b tweaking 2019-09-27 16:38:11 -05:00
Chao Liu
14315b72f3 tweaking 2019-09-27 15:24:27 -05:00
Chao Liu
ebe38f3d48 debugging 2019-09-27 11:31:01 -05:00
Chao Liu
9b280cc50d remove dead code 2019-09-27 02:00:59 -05:00
Chao Liu
98a2cfcc84 nvidia build 2019-09-27 00:15:05 -05:00
Chao Liu
00089cd6e5 clean up 2019-09-26 21:39:28 -05:00
Chao Liu
b12bbceebc clean up 2019-09-26 14:59:19 -05:00
Chao Liu
51a9fa1dbd removing dependency on old tensor descriptor 2019-09-26 11:49:05 -05:00
Chao Liu
0f52c4c0e4 added type conversion in threadwise and blockwise copy 2019-09-26 00:00:25 -05:00
Chao Liu
b3d4595f5a added type conversion in threadwise and blockwise copy 2019-09-25 23:38:26 -05:00
Chao Liu
3cb2a7d09f removing old implementation of tensor descriptor 2019-09-25 22:43:34 -05:00
Chao Liu
39d92e7dfd removing old implementation of tensor descriptor 2019-09-25 22:24:06 -05:00
Chao Liu
012b525377 clean up 2019-09-25 03:28:53 -05:00
Chao Liu
e1ae8f18f7 added GetLinearDimensionMask 2019-09-25 02:52:41 -05:00
Chao Liu
4f4aba4872 adding GetLinearDimensionMask() 2019-09-24 23:59:47 -05:00
Chao Liu
545d930568 refactor 2019-09-24 18:06:05 -05:00
Chao Liu
37f4e2b6d8 nvidia build 2019-09-22 03:23:19 -05:00
Chao Liu
6c2c50b020 done: explicitly separate offset component into compile-time, block-invariant and per-thread components. Experimenting 2019-09-22 03:17:41 -05:00
Chao Liu
51884fc214 WIP: explicitly separate offset component into compile-time, block-invariant and per-thread components 2019-09-21 22:53:03 -05:00
Chao Liu
740da00aa2 refactor 2019-09-20 21:45:20 -05:00
Chao Liu
184c6e7d37 nvidia build 2019-09-20 21:45:03 -05:00
Chao Liu
f00c138145 adding logic to judge linear dimension 2019-09-20 20:43:13 -05:00
Chao Liu
bf7e7d62a8 refactor 2019-09-19 23:44:23 -05:00
Chao Liu
b6e1c52a80 use buffer_load buffer_store intrinsic 2019-09-19 15:39:07 -05:00
Chao Liu
8afbb10d18 reduce some register usage in index 2019-09-18 16:42:45 -05:00
Chao Liu
94bb1b4835 refactor 2019-09-18 16:08:24 -05:00
Chao Liu
86cc678f18 add global_load and buffer_load inline asm 2019-09-18 15:41:55 -05:00
Chao Liu
5b7a18c506 experimenting global and buffer load/store 2019-09-18 02:05:42 -05:00
Chao Liu
c7a6545ec4 experimenting global and buffer load/store 2019-09-18 01:37:28 -05:00
Chao Liu
9f46cdf5fa experimenting global and buffer load/store 2019-09-18 00:15:57 -05:00
Chao Liu
f58bf38445 enable hip compiler flag: -amdgpu-enable-global-sgpr-addr 2019-09-17 17:34:39 -05:00
Chao Liu
126cae0c9b bug fix 2019-09-17 15:02:12 -05:00
Chao Liu
e1a67b693e refactor 2019-09-17 11:19:15 -05:00
Chao Liu
f7be86b9e4 refactor 2019-09-16 22:47:55 -05:00
Chao Liu
d707993933 bug fix 2019-09-15 20:57:07 -05:00
Chao Liu
69fea593ec amd build 2019-09-15 17:55:46 -05:00
Chao Liu
940949d9d5 add lds doble buffer to nchw padded v4r1 and v4r4 2019-09-15 16:59:54 -05:00
Chao Liu
bf97542846 add lds doble buffer to nchw padded v4r1 and v4r4 2019-09-15 16:58:16 -05:00