Commit Graph

72 Commits

Author SHA1 Message Date
Chao Liu
3bc034ec9e tweaking
[ROCm/composable_kernel commit: 012d3a071b]
2019-09-27 16:38:11 -05:00
Chao Liu
28d6a9834a tweaking
[ROCm/composable_kernel commit: 14315b72f3]
2019-09-27 15:24:27 -05:00
Chao Liu
8df18296f5 debugging
[ROCm/composable_kernel commit: ebe38f3d48]
2019-09-27 11:31:01 -05:00
Chao Liu
a87fa81015 remove dead code
[ROCm/composable_kernel commit: 9b280cc50d]
2019-09-27 02:00:59 -05:00
Chao Liu
1109bb05d0 clean up
[ROCm/composable_kernel commit: 00089cd6e5]
2019-09-26 21:39:28 -05:00
Chao Liu
c5e4a623f9 clean up
[ROCm/composable_kernel commit: b12bbceebc]
2019-09-26 14:59:19 -05:00
Chao Liu
bd8f263b70 removing dependency on old tensor descriptor
[ROCm/composable_kernel commit: 51a9fa1dbd]
2019-09-26 11:49:05 -05:00
Chao Liu
313eb881cd removing old implementation of tensor descriptor
[ROCm/composable_kernel commit: 39d92e7dfd]
2019-09-25 22:24:06 -05:00
Chao Liu
66f7a03ee4 clean up
[ROCm/composable_kernel commit: 012b525377]
2019-09-25 03:28:53 -05:00
Chao Liu
ca661e1f52 refactor
[ROCm/composable_kernel commit: 545d930568]
2019-09-24 18:06:05 -05:00
Chao Liu
c3a1be3865 WIP: explicitly separate offset component into compile-time, block-invariant and per-thread components
[ROCm/composable_kernel commit: 51884fc214]
2019-09-21 22:53:03 -05:00
Chao Liu
b9722daae4 refactor
[ROCm/composable_kernel commit: bf7e7d62a8]
2019-09-19 23:44:23 -05:00
Chao Liu
5d1fb91f12 refactor
[ROCm/composable_kernel commit: 94bb1b4835]
2019-09-18 16:08:24 -05:00
Chao Liu
67f93cb50c add global_load and buffer_load inline asm
[ROCm/composable_kernel commit: 86cc678f18]
2019-09-18 15:41:55 -05:00
Chao Liu
6c5f82174b experimenting global and buffer load/store
[ROCm/composable_kernel commit: 5b7a18c506]
2019-09-18 02:05:42 -05:00
Chao Liu
741a647405 experimenting global and buffer load/store
[ROCm/composable_kernel commit: c7a6545ec4]
2019-09-18 01:37:28 -05:00
Chao Liu
03740b5c4a experimenting global and buffer load/store
[ROCm/composable_kernel commit: 9f46cdf5fa]
2019-09-18 00:15:57 -05:00
Chao Liu
0342a42f47 enable hip compiler flag: -amdgpu-enable-global-sgpr-addr
[ROCm/composable_kernel commit: f58bf38445]
2019-09-17 17:34:39 -05:00
Chao Liu
e9653e9ecc bug fix
[ROCm/composable_kernel commit: 126cae0c9b]
2019-09-17 15:02:12 -05:00
Chao Liu
2b26fb76ca refactor
[ROCm/composable_kernel commit: f7be86b9e4]
2019-09-16 22:47:55 -05:00
Chao Liu
ff99bc281e bug fix
[ROCm/composable_kernel commit: d707993933]
2019-09-15 20:57:07 -05:00
Chao Liu
d89b9c2e08 amd build
[ROCm/composable_kernel commit: 69fea593ec]
2019-09-15 17:55:46 -05:00
Chao Liu
7ff1c75c9d add lds doble buffer to nchw padded v4r1 and v4r4
[ROCm/composable_kernel commit: 940949d9d5]
2019-09-15 16:59:54 -05:00
Chao Liu
a4edaf2ae3 add lds doble buffer to nchw padded v4r1 and v4r4
[ROCm/composable_kernel commit: bf97542846]
2019-09-15 16:58:16 -05:00
Chao Liu
e6285a3b55 initial implementation for nchw v4r4 padding
[ROCm/composable_kernel commit: 2c93b3057d]
2019-09-15 16:31:54 -05:00
Chao Liu
89dc09f929 clean up
[ROCm/composable_kernel commit: 53094f7fae]
2019-09-15 12:13:58 -05:00
Chao Liu
93c96db8c8 initial padding support for nchw
[ROCm/composable_kernel commit: d4878d99f9]
2019-09-13 23:30:48 -05:00
Chao Liu
79e13d9ded clean up
[ROCm/composable_kernel commit: bd7a230006]
2019-09-12 14:55:46 -05:00
Chao Liu
bea9be7a3c padding for chwn is functional
[ROCm/composable_kernel commit: 1f70524471]
2019-09-12 01:12:08 -05:00
Chao Liu
9a8d9e0e40 enabling padding for chwn format
[ROCm/composable_kernel commit: 724e984bff]
2019-09-11 01:13:13 -05:00
Chao Liu
17564ecfec adding merge transform
[ROCm/composable_kernel commit: ca42e9101d]
2019-09-10 01:53:49 -05:00
Chao Liu
399be319a2 more utility code
[ROCm/composable_kernel commit: 7a7fe16086]
2019-09-09 00:29:33 -05:00
Chao Liu
b0f3708397 added tuple
[ROCm/composable_kernel commit: 625838def0]
2019-09-06 18:07:56 -05:00
Chao Liu
4c17a5a102 clean
[ROCm/composable_kernel commit: cb6475c77d]
2019-08-23 09:59:23 -05:00
Chao Liu
a74c7c8e8a adding tensor_view
[ROCm/composable_kernel commit: 238d58c2f5]
2019-08-20 17:29:54 -05:00
Chao Liu
eb6a36d393 Merge remote-tracking branch 'origin/master' into add_padding
[ROCm/composable_kernel commit: 86ceded98b]
2019-08-15 13:48:45 -05:00
Chao Liu
c8e088723f clean up
[ROCm/composable_kernel commit: 0979fb4af9]
2019-08-15 13:21:51 -05:00
Chao Liu
3b234309c9 adding padding to implicit gemm v1r3
[ROCm/composable_kernel commit: 4fb81e008c]
2019-08-14 10:55:34 -05:00
Chao Liu
9c00690649 clean up
[ROCm/composable_kernel commit: 740149fcf1]
2019-08-13 17:26:00 -05:00
Chao Liu
15a2fc9029 add back some code
[ROCm/composable_kernel commit: 40836ab926]
2019-08-13 12:21:38 -05:00
Chao Liu
18bc57cd93 clean up
[ROCm/composable_kernel commit: 8bdaba51f8]
2019-08-13 00:37:23 -05:00
Chao Liu
301348f208 cleaning up
[ROCm/composable_kernel commit: 1c4ef23cff]
2019-08-09 22:48:28 -05:00
Chao Liu
807d22de0b tweak on amd
[ROCm/composable_kernel commit: 4908fe3fdc]
2019-08-08 12:14:06 -05:00
Chao Liu
0b8efbb3e9 clean up
[ROCm/composable_kernel commit: 701b7341f0]
2019-08-07 19:25:54 -05:00
Chao Liu
139b043d07 use ford/for instead of static_ford/static_for in threadwise copy, somehow register spill is greatly reduced on AMD
[ROCm/composable_kernel commit: bc9ea646f8]
2019-08-07 19:09:13 -05:00
Chao Liu
8994adaca4 reworked ThreadwiseGenericTensorSliceCopy_v1
[ROCm/composable_kernel commit: 1b3c2e4035]
2019-08-07 00:52:13 -05:00
Chao Liu
fb1953ed86 reimplement threadwise copy
[ROCm/composable_kernel commit: fdcfae3a62]
2019-08-06 17:41:58 -05:00
Chao Liu
0319e69e81 tweak
[ROCm/composable_kernel commit: adc1008836]
2019-08-03 15:05:25 -05:00
Chao Liu
783cc11cef tweak
[ROCm/composable_kernel commit: 4a1e97cf86]
2019-08-03 14:33:39 -05:00
Chao Liu
be8343bb7e added implicit gemm v4r4 and double buffer
[ROCm/composable_kernel commit: c2d246696f]
2019-08-03 00:19:19 -05:00