ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-14 02:02:46 +00:00
Code Issues Packages Projects Releases Wiki Activity
453 Commits 742 Branches 38 Tags
abda4c0cab93c74762c157a4d15e8e04777d75be
Commit Graph

453 Commits

This Branch
This Branch
All Branches
Author SHA1 Message Date
Qianfeng
abda4c0cab [MIOpen Downstream] Fix Reduction Kernel (#34) 
* Tiny fix in using data type template parameters in blockwise and direct_threadwise kernel

* Fix with regard to implementing GetZeroVal() in both kernel and host

* Avoid convert to compType from dstDataType before writting the output value

* Add half_t support to NumericLimits and make constexpr GetZeroVal() of binary operator

* Add CONSTANT decorator for descriptor read buffer

* Use get_thread_local_1d_id() for thread local Id

* Rename GetZeroVal() to GetReductionZeroVal() in the kernels

* Remove constexpr from initialized zeroVal and tiny fix in reduction_operator.hpp

* Occasional tiny simplification and update in the kernel files

* Update to re-order tensor dimensions on the host, split second_call kernel wrapper files and simplify reduce_all kernel wrappers

* Update to remove OpenCL tidy checking failures

* Update for better readability

* Remove unused codes and not-needed template parameters in the kernel wrappers

Co-authored-by: Chao Liu <chao.liu2@amd.com>

[ROCm/composable_kernel commit: b2dc55f82c]
 2021-10-06 14:43:17 -05:00
Chao Liu
0e75841071 Tweak GEMM kernel (#38) 
* add parameters

* tweak gemm

* tweak

* update conv

* update script

* adding bwd 1x1

* update script

* adding 1x1 bwd

* debugging bwd 1x1 failure

* update script

* update script

* test

* test v100

* clean up

[ROCm/composable_kernel commit: b3e8d57d51]
 2021-10-06 11:12:36 -05:00
zjing14
ad110e92ba Add VectorType support into StaticBuffer (#27) 
* init StaticBufferV2

* clean

* adopt old output stage for staticBufferV2

* clean

* remove hack

* clean

* clean

* clean code

* move c_buffer alloc into blockwise gemm

* add adaptors for m/n_thread_data_on_grid

* adjust blockwise_gemm_xdlops

* reorder ops in GEMM hot loop

Co-authored-by: Chao Liu <chao.liu2@amd.com>

[ROCm/composable_kernel commit: 846f462bd4]
 2021-10-06 10:13:52 -05:00
Chao Liu
c732ff2164 Merge remote-tracking branch 'origin/develop' into miopen_downstream-dynamic_reduction_pr 
[ROCm/composable_kernel commit: b725e3fc84]
 2021-09-21 11:55:26 -05:00
Chao Liu
b047701d37 :Merge remote-tracking branch 'origin/develop' into CK_upstream 
[ROCm/composable_kernel commit: df0d68106e]
 2021-09-20 20:44:01 -05:00
Chao Liu
5141373604 Add a version of Merge transform that use integerdivision and mod (#25) 
* add Merg_v3_division_mod

* refactor

[ROCm/composable_kernel commit: f3acd2510b]
 2021-09-05 12:57:57 -05:00
Chao Liu
115f77e17a GEMM driver and kernel (#29) 
* add gemm driver

* tweak

* add gemm kernel: mk_kn_mn and km_kn_mn

* tweak

* add GEMM km_nk_mn

* fix comment

[ROCm/composable_kernel commit: 19613902b5]
 2021-09-05 12:41:28 -05:00
ltqin
79b671c5dd Backward weight v4r4r2 with xdlops (#18) 
* start

* modify transformat

* modify device convolutiion

* modify host

* added host conv bwd and wrw

* remove bwd, seperate wrw

* clean

* hacall k to zero

* out log

* fixed

* fixed

* change to (out in wei)

* input hack

* hack to out

* format

* fix by comments

* change wei hacks(wei transform has not merge)

* fix program once issue

* fix review comment

* fix vector load issue

* tweak

Co-authored-by: ltqin <letaoqin@amd.com>
Co-authored-by: Jing Zhang <jizhan@amd.com>
Co-authored-by: Chao Liu <chao.liu2@amd.com>

[ROCm/composable_kernel commit: 627d8ef35a]
 2021-08-30 22:49:17 -05:00
Chao Liu
0dc65eae46 Misc fixes (#24) 
* use cast_pointer_to_generic_address_space() in v6r1 kernel wrapper, DynamcBuffer and buffer_load take customized invalid-element-value, add buffer_load/store for fp64

* use remove_cvref_t

[ROCm/composable_kernel commit: 10bb811060]
 2021-08-26 20:05:19 -05:00
Qianfeng
b315c39b11 [SWDEV-281541][MSRCHA-100] Implementation of Dynamic Generic Reduction (#1108) 
* add solver ConvIgemmFwdV6r1DlopsNchwKcyxNkhw; rename static ck source files

* make inner product compatible on gfx900

* Update src/include/miopen/solver/ck_utility_common.hpp

* compiler parameter use stream

* use int instead of index_t in kernel wrapper

* DynamicBuffer, StaticBuffer, amd_buffer_load support customized value for invalid element

* Add dynamic generic reduction kernel layer (kernel wrappers, kernel implementations and utilities)

* Some updates to dynamic composable kernel facility for the need of dynamic generic reduction

* Update to generic reduction C++ host interface layer to support dynamic generic reduction

* Update to remove tidy complaints in host interface layer

* Change the unary operator form from void op(T &x) to T op(T x)

* Update to pass single workspace pointer for all kernels (fix for OpenCL backend)

* Use cppcheck-suppress to prevent some strange warnings

* Re-use operator [] and () for DynamicBuffer and update to depending codes

* Remove useless codes in first call threadwise/warpwise/blockwise kernel wrappers

* [performance] Remove un-needed local buffer initialization

Co-authored-by: Chao Liu <chao.liu2@amd.com>
Co-authored-by: JD <Jehandad.Khan@amd.com>

[ROCm/composable_kernel commit: 9e80cdceb7]
 2021-08-26 18:04:55 -07:00
zjing14
346651889b GlobalAtomicAdd for fp32/int32 (#23) 
* add f32/i32 atomicAdd support into dynamicBuffer, and enable it in v1r3

* fixed

* fixed

* update comment

Co-authored-by: Chao Liu <chao.liu2@amd.com>

[ROCm/composable_kernel commit: a7a758d8ce]
 2021-08-25 10:55:55 -05:00
zjing14
6d56ee7d00 Xdlops refactor fix (#22) 
* added constexpr ahead of adptor; clean unused driver; rename M/NPerWave to M/NPerXDL

* fixed bwd

* fixed comment

[ROCm/composable_kernel commit: 9d3f634a3c]
 2021-08-23 11:22:10 -05:00
Chao Liu
818e5cbfe7 magic division use __umulhi() (#19) 
[ROCm/composable_kernel commit: c6f26bb480]
 2021-08-23 10:40:27 -05:00
Chao Liu
ee428d2d6f Composable kernel init integration v3 (#1097) 
* Squashed 'src/composable_kernel/' content from commit f6edda611

git-subtree-dir: src/composable_kernel
git-subtree-split: f6edda6119

* add solver ConvIgemmFwdV6r1DlopsNchwKcyxNkhw; rename static ck source files

* Squashed 'src/composable_kernel/' changes from f6edda611..5781adf5c

5781adf5c Update develop (#5) (#6)
97e6d514f Merge pull request #4 from ROCmSoftwarePlatform/separate_online_compile
7b1ec41e5 refactor
49c33aaea refactor
54b3e73d1 rename

git-subtree-dir: src/composable_kernel
git-subtree-split: 5781adf5cf

* fix

* refactor

* remove online compilation from CK

* refactor

* fix

* add ctest

* add c-style pointer cast

* vector/scalar pointer cast use c-style pointer cast instead of reinterpret_cast

* fix clang warning suppression

* tidy

* suppress cppcheck

* fix enum issue

* revert chagnes to hip build

* fix kernel filename

* update CK build script

* rename

* rename

* make innner product compatiable on gfx900

* Update src/include/miopen/solver/ck_utility_common.hpp

Co-authored-by: JD <Jehandad.Khan@amd.com>

* compiler parameter use stream

* use int instead of index_t in kernel wrapper

* DynamicBuffer, StaticBuffer, amd_buffer_load support customized value for invalid element

* refactor

* refactor

* change cmakelist

* change ck common utility

* fix

Co-authored-by: JD <Jehandad.Khan@amd.com>

[ROCm/composable_kernel commit: 6fe3627a9e]
 2021-08-19 10:55:03 -05:00
zjing14
3cacb0c037 refactor dynamic xdlops iGemm (#13) 
* xdlops refactor

* fixed commnt

* clean xdlops_gemm

* add make c into xldops-gemm

* change mfma_info

* refactor xdlops, hide c desc

* clean

* clean

* clean

* apply hacks changes to v4r4r4_nhwc

* rename hacks and use single stage adapter

* enable fp16 mfma

[ROCm/composable_kernel commit: a2ad6d3531]
 2021-08-19 09:54:10 -05:00
zjing14
13dab08075 Added host_conv_wrw for verification (#15) 
* added host conv wrw

[ROCm/composable_kernel commit: ba6f79a75e]
 2021-08-19 01:00:41 -05:00
Chao Liu
dd3e30e224 Merge pull request #14 from ROCmSoftwarePlatform/miopen_downstream_init_integration 
MIOpen Downstream: Initial integration 2nd PR

[ROCm/composable_kernel commit: b62bf8c3f8]
 2021-08-16 16:39:40 -05:00
Chao Liu
a364868bcf Merge pull request #8 from ROCmSoftwarePlatform/miopen_downstream_init_integration 
[ROCm/composable_kernel commit: ccc4a1d365]
 2021-08-16 16:28:53 -05:00
Chao Liu
bf335705ef refactor 
[ROCm/composable_kernel commit: 67ad47e7c1]
 2021-08-16 21:01:33 +00:00
Chao Liu
04d90a65c5 refactor 
[ROCm/composable_kernel commit: 16effa767c]
 2021-08-16 20:36:47 +00:00
Chao Liu
a222af4530 DynamicBuffer, StaticBuffer, amd_buffer_load support customized value for invalid element 
[ROCm/composable_kernel commit: a91b68dfcb]
 2021-08-13 23:40:19 +00:00
Chao Liu
e145aa0e29 use int instead of index_t in kernel wrapper 
[ROCm/composable_kernel commit: 2cbabbba54]
 2021-08-13 20:55:39 +00:00
Chao Liu
d49b0e5239 compiler parameter use stream 
[ROCm/composable_kernel commit: 0834bc7635]
 2021-08-13 01:05:14 +00:00
Chao Liu
c5ea28dc43 make innner product compatiable on gfx900 
[ROCm/composable_kernel commit: f2ac7832c6]
 2021-08-11 09:42:53 -05:00
Chao Liu
0bbcd05e4e rename 
[ROCm/composable_kernel commit: 4e57b30a6a]
 2021-08-11 00:08:42 +00:00
Chao Liu
c5a4edb9e8 rename 
[ROCm/composable_kernel commit: c03045ce2d]
 2021-08-10 23:45:36 +00:00
Chao Liu
400efa2b88 update CK build script 
[ROCm/composable_kernel commit: b2589957f3]
 2021-08-10 22:19:13 +00:00
Chao Liu
a0df0eb029 fix kernel filename 
[ROCm/composable_kernel commit: 2c48039d0e]
 2021-08-10 22:15:23 +00:00
Chao Liu
7d53faec4e fix enum issue 
[ROCm/composable_kernel commit: d626dccc95]
 2021-08-10 20:55:13 +00:00
Chao Liu
2e7587d6db tidy 
[ROCm/composable_kernel commit: 643ebd4f3e]
 2021-08-10 07:07:11 +00:00
Chao Liu
a15f7c025f fix clang warning suppression 
[ROCm/composable_kernel commit: ddd49ec9e7]
 2021-08-10 06:20:24 +00:00
Chao Liu
ab465fca4c vector/scalar pointer cast use c-style pointer cast instead of reinterpret_cast 
[ROCm/composable_kernel commit: 4f566c6221]
 2021-08-10 05:55:20 +00:00
Chao Liu
17b084ce2c add c-style pointer cast 
[ROCm/composable_kernel commit: 172036d728]
 2021-08-10 00:01:52 -05:00
Chao Liu
72c6d11a3b tidy 
[ROCm/composable_kernel commit: 76f3131939]
 2021-08-09 18:49:59 -05:00
Chao Liu
7b306ba01c tidy 
[ROCm/composable_kernel commit: d18428901e]
 2021-08-09 18:20:02 -05:00
Chao Liu
eb778cb91b tidy 
[ROCm/composable_kernel commit: f885c131d8]
 2021-08-09 22:13:47 +00:00
Chao Liu
922b72948d tidy 
[ROCm/composable_kernel commit: 80120f0a0c]
 2021-08-09 21:10:09 +00:00
Chao Liu
a913d48b07 tidy 
[ROCm/composable_kernel commit: c3efeb5e20]
 2021-08-09 19:32:07 +00:00
Chao Liu
590dde14c8 tidy 
[ROCm/composable_kernel commit: 56fc0842b3]
 2021-08-09 19:27:49 +00:00
Chao Liu
7885261dc6 tidy 
[ROCm/composable_kernel commit: 54fba515b3]
 2021-08-09 17:33:32 +00:00
Chao Liu
193f7cbbcf tidy 
[ROCm/composable_kernel commit: e62bae7a4a]
 2021-08-09 15:11:35 +00:00
Chao Liu
71a28fc501 add tidy 
[ROCm/composable_kernel commit: 24c8728942]
 2021-08-08 17:41:54 +00:00
Chao Liu
c2922caeb1 fix 
[ROCm/composable_kernel commit: 61487e0a00]
 2021-08-07 02:31:19 +00:00
Chao Liu
b0ab2054a9 remove online compilation from CK 
[ROCm/composable_kernel commit: ae98b52ad8]
 2021-08-07 00:51:05 +00:00
Chao Liu
23239fa618 refactor 
[ROCm/composable_kernel commit: cb95421311]
 2021-08-06 22:17:51 +00:00
Chao Liu
3bbe4dae99 Merge commit '437cc595c6e206dfebb118985b5171bbc1e29eab' into composable_kernel_init_integration_v3 
[ROCm/composable_kernel commit: 73ca970155]
 2021-08-06 21:32:27 +00:00
Chao Liu
51d5010798 Merge pull request #7 from ROCmSoftwarePlatform/master 
Master

[ROCm/composable_kernel commit: 3b86646115]
 2021-08-06 16:29:58 -05:00
Chao Liu
98d5803d26 Update develop (#5) 
* refactor

[ROCm/composable_kernel commit: d09ea4f4e5]
 2021-08-06 16:11:15 -05:00
Chao Liu
37a9932cfc refactor 
[ROCm/composable_kernel commit: 7b1ec41e5b]
 2021-08-06 20:50:01 +00:00
Chao Liu
101d8ecf21 refactor 
[ROCm/composable_kernel commit: 49c33aaea7]
 2021-08-06 19:59:53 +00:00