Commit Graph

335 Commits

Author SHA1 Message Date
Chao Liu
3bd3443225 adding logic to judge linear dimension
[ROCm/composable_kernel commit: f00c138145]
2019-09-20 20:43:13 -05:00
Chao Liu
b9722daae4 refactor
[ROCm/composable_kernel commit: bf7e7d62a8]
2019-09-19 23:44:23 -05:00
Chao Liu
056e48a3ac use buffer_load buffer_store intrinsic
[ROCm/composable_kernel commit: b6e1c52a80]
2019-09-19 15:39:07 -05:00
Chao Liu
c1f7320764 reduce some register usage in index
[ROCm/composable_kernel commit: 8afbb10d18]
2019-09-18 16:42:45 -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
eb1bb96a52 refactor
[ROCm/composable_kernel commit: e1a67b693e]
2019-09-17 11:19:15 -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
d0a36e9a40 adding dimension transformation
[ROCm/composable_kernel commit: 12da8154c8]
2019-09-05 00:20:05 -05:00
Chao Liu
6166bf61c2 adding dimension tranformation
[ROCm/composable_kernel commit: 0c05f4279f]
2019-09-05 00:19:06 -05:00
Chao Liu
572dd9ae5c adding dimension transformation
[ROCm/composable_kernel commit: bd44e6390d]
2019-09-02 00:21:00 -05:00
Chao Liu
4c17a5a102 clean
[ROCm/composable_kernel commit: cb6475c77d]
2019-08-23 09:59:23 -05:00
Chao Liu
e1b3f164cd Merge remote-tracking branch 'origin/master' into add_tensor_view
[ROCm/composable_kernel commit: 6ff3fe5d05]
2019-08-21 14:46:16 -05:00
Chao Liu
dd6fef3b9e add script for doing Jack's ISA injection hack
[ROCm/composable_kernel commit: 0c83df668f]
2019-08-21 14:29:13 -05:00
Chao Liu
a74c7c8e8a adding tensor_view
[ROCm/composable_kernel commit: 238d58c2f5]
2019-08-20 17:29:54 -05:00
Chao Liu
c0d154b07d bug fix: BlockwiseGenericTensorSliceCopy_v2::MoveDstSlicingWindow
[ROCm/composable_kernel commit: 08bf57b01c]
2019-08-15 15:12:13 -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
d54e146f28 clean up
[ROCm/composable_kernel commit: fab2f10a55]
2019-08-12 15:48:35 -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
7b5f9bebc6 added ThreadwiseGenericTensorSliceCopy_v2r1
[ROCm/composable_kernel commit: a9b2b1dcd7]
2019-08-08 02:42:52 -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
c9a4be4b22 bug fix in ford, forgot to reorder lengths
[ROCm/composable_kernel commit: 5636576f9b]
2019-08-07 18:27:10 -05:00
Chao Liu
5716cecadc adding ThreadwiseGenericTensorSliceCopy_v1r2
[ROCm/composable_kernel commit: 9d99a58072]
2019-08-07 16:51:14 -05:00
Chao Liu
8994adaca4 reworked ThreadwiseGenericTensorSliceCopy_v1
[ROCm/composable_kernel commit: 1b3c2e4035]
2019-08-07 00:52:13 -05:00
Chao Liu
56ad0dcd3f add looping Orders into ford and static_ford
[ROCm/composable_kernel commit: 41cdde99e5]
2019-08-06 20:23:11 -05:00