ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-15 18:42:06 +00:00
Code Issues Packages Projects Releases Wiki Activity
370 Commits 745 Branches 38 Tags
ceffe74a0a567a8ea129e47bac9fbd6ef49d0444
Commit Graph

370 Commits

This Branch
This Branch
All Branches
Author SHA1 Message Date
Chao Liu
ceffe74a0a Improve buffer address for out of bound check (#21) 
* Use buffer load built-in OOB check. buffer size is limited to 2GB.
* buffer APIs use combined wave and thread offset
* use uint32_t for addr shift in buffer addressing

[ROCm/composable_kernel commit: ac62d13ecd]
 2020-07-29 18:04:09 -05:00
Chao Liu
a3c89131fa Code clean up (#20) 
* tuning para,

* testing on v100

* add fp16

* remove deprecated tensor descriptor

* sync with miopen

* update build script

Co-authored-by: Jing Zhang <jizhan@amd.com>

[ROCm/composable_kernel commit: 5c7cec1115]
 2020-06-23 20:31:27 -05:00
Chao Liu
82d8fb8494 MIOpen integration (#15) 
* renaming

[ROCm/composable_kernel commit: 7d09790a0a]
 2020-02-18 10:42:18 -06:00
Chao Liu
ed1eafcec8 MIopen integration (#13) 
* update for miopen integration: cosmetic refactor


[ROCm/composable_kernel commit: 1a66e35b6f]
 2020-02-17 09:53:20 -06:00
Chao Liu
bd24dfbea7 Update for recent MIOpen integration (#11) 
* update for MIOpen integration


[ROCm/composable_kernel commit: 3406a1148a]
 2020-01-27 15:29:33 -06:00
Chao Liu
7c9100b53f 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

[ROCm/composable_kernel commit: c5da0377fb]
 2020-01-20 10:20:03 -06:00
Chao Liu
24f7d66609 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

[ROCm/composable_kernel commit: e2b4c5b469]
 2019-12-05 12:36:36 -06:00
Chao Liu
b075179dd9 fixed faulty padding API calls (#8) 
[ROCm/composable_kernel commit: 19a93dac05]
 2019-12-03 01:46:44 -06:00
Chao Liu
4414e495ed 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


[ROCm/composable_kernel commit: 8f5f64960e]
 2019-12-03 01:16:12 -06:00
Chao Liu
809e8239e1 remove dead file (#6) 
[ROCm/composable_kernel commit: 31ded4ac4b]
 2019-11-04 17:13:38 -06:00
Chao Liu
b58da02f74 MIOpen integration: recent bug fixes from MIOpen (#5) 
[ROCm/composable_kernel commit: 562e1e2767]
 2019-11-04 16:51:12 -06:00
Chao Liu
a9e6c3340c Refactor for MIOpen integration (#4) 
Refactor, so can bring multi-index transformation and padding support into MIOpen

[ROCm/composable_kernel commit: 52c3fe05be]
 2019-10-11 11:37:31 -05:00
Chao Liu
2327a08c21 Merge pull request #3 from asroy/clean_up 
enable type conversion in ThreadwiseGenericTensorSliceCopy_v2r1 and BlockwiseGenericTensorSliceCopy_v2

[ROCm/composable_kernel commit: 9aaeacc82b]
 2019-09-30 15:14:02 -05:00
Chao Liu
0c41efb629 enable type conversion in blockwise copy v2 and threadwise copy v2r1 
[ROCm/composable_kernel commit: cf21818455]
 2019-09-30 15:11:05 -05:00
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
38b3aeafb6 nvidia build 
[ROCm/composable_kernel commit: 98a2cfcc84]
 2019-09-27 00:15:05 -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
eb0a575ad1 added type conversion in threadwise and blockwise copy 
[ROCm/composable_kernel commit: 0f52c4c0e4]
 2019-09-26 00:00:25 -05:00
Chao Liu
4b54b1f190 added type conversion in threadwise and blockwise copy 
[ROCm/composable_kernel commit: b3d4595f5a]
 2019-09-25 23:38:26 -05:00
Chao Liu
2c7fc0bc28 removing old implementation of tensor descriptor 
[ROCm/composable_kernel commit: 3cb2a7d09f]
 2019-09-25 22:43:34 -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
9695b9f548 added GetLinearDimensionMask 
[ROCm/composable_kernel commit: e1ae8f18f7]
 2019-09-25 02:52:41 -05:00
Chao Liu
1019fcf64d adding GetLinearDimensionMask() 
[ROCm/composable_kernel commit: 4f4aba4872]
 2019-09-24 23:59:47 -05:00
Chao Liu
ca661e1f52 refactor 
[ROCm/composable_kernel commit: 545d930568]
 2019-09-24 18:06:05 -05:00
Chao Liu
9338f82689 nvidia build 
[ROCm/composable_kernel commit: 37f4e2b6d8]
 2019-09-22 03:23:19 -05:00
Chao Liu
b60cecbbc8 done: explicitly separate offset component into compile-time, block-invariant and per-thread components. Experimenting 
[ROCm/composable_kernel commit: 6c2c50b020]
 2019-09-22 03:17:41 -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
f961bfabf7 refactor 
[ROCm/composable_kernel commit: 740da00aa2]
 2019-09-20 21:45:20 -05:00
Chao Liu
18b23028f0 nvidia build 
[ROCm/composable_kernel commit: 184c6e7d37]
 2019-09-20 21:45:03 -05:00
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