ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-11 17:51:40 +00:00
Code Issues Packages Projects Releases Wiki Activity
1,961 Commits 778 Branches 40 Tags
streamk_debug
Commit Graph

1961 Commits

This Branch
This Branch
All Branches
Author SHA1 Message Date
Chao Liu
b3e8d57d51 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
 2021-10-06 11:12:36 -05:00
zjing14
846f462bd4 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>
 2021-10-06 10:13:52 -05:00
Qianfeng
dfb80c4e39 [Enhancements] Several bugfixes and refactoring of dynamic generic reduction (#1156) 
* 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

* tidy

* add tidy

* tidy

* tidy

* tidy

* tidy

* tidy

* tidy

* tidy

* tidy

* tidy

* 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

* Squashed 'src/composable_kernel/' changes from 5781adf5c..31b403526

31b403526 Merge pull request #16 from ROCmSoftwarePlatform/develop
b62bf8c3f Merge pull request #14 from ROCmSoftwarePlatform/miopen_downstream_init_integration
ccc4a1d36 Merge pull request #8 from ROCmSoftwarePlatform/miopen_downstream_init_integration
67ad47e7c refactor
16effa767 refactor
a91b68dfc DynamicBuffer, StaticBuffer, amd_buffer_load support customized value for invalid element
2cbabbba5 use int instead of index_t in kernel wrapper
0834bc763 compiler parameter use stream
f2ac7832c make innner product compatiable on gfx900
4e57b30a6 rename
c03045ce2 rename
b2589957f update CK build script
2c48039d0 fix kernel filename
d626dccc9 fix enum issue
643ebd4f3 tidy
ddd49ec9e fix clang warning suppression
4f566c622 vector/scalar pointer cast use c-style pointer cast instead of reinterpret_cast
172036d72 add c-style pointer cast
76f313193 tidy
d18428901 tidy
f885c131d tidy
80120f0a0 tidy
c3efeb5e2 tidy
56fc0842b tidy
54fba515b tidy
e62bae7a4 tidy
24c872894 add tidy
61487e0a0 fix
ae98b52ad remove online compilation from CK
cb9542131 refactor
73ca97015 Merge commit '437cc595c6e206dfebb118985b5171bbc1e29eab' into composable_kernel_init_integration_v3
3b8664611 Merge pull request #7 from ROCmSoftwarePlatform/master
d09ea4f4e Update develop (#5)
3d32ae940 add solver ConvIgemmFwdV6r1DlopsNchwKcyxNkhw; rename static ck source files

git-subtree-dir: src/composable_kernel
git-subtree-split: 31b403526e

* 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 in src/reducetensor.cpp for consistent IDs passing to the kernel

* 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

* Small updates in src/reducetensor.cpp

* 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>
Co-authored-by: JD <Jehandad.Khan@amd.com>
 2021-09-29 08:12:11 -07:00
Jun Liu
8557901d02 Merge pull request #1165 from ROCmSoftwarePlatform/develop 
Merge develop into CK_upstream (Please don't squash when merging)
 2021-09-21 15:52:12 -07:00
Chao Liu
f305bebdc3 Merge pull request #31 from ROCmSoftwarePlatform/miopen_downstream-dynamic_reduction_pr 
[MIOpen Downstream] Dynamic Reduction PR
 2021-09-21 11:59:23 -05:00
Chao Liu
b725e3fc84 Merge remote-tracking branch 'origin/develop' into miopen_downstream-dynamic_reduction_pr 2021-09-21 11:55:26 -05:00
Chao Liu
df0d68106e :Merge remote-tracking branch 'origin/develop' into CK_upstream 2021-09-20 20:44:01 -05:00
Chao Liu
f3acd2510b Add a version of Merge transform that use integerdivision and mod (#25) 
* add Merg_v3_division_mod

* refactor
 2021-09-05 12:57:57 -05:00
Chao Liu
19613902b5 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
 2021-09-05 12:41:28 -05:00
ltqin
627d8ef35a 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>
 2021-08-30 22:49:17 -05:00
Chao Liu
10bb811060 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
 2021-08-26 20:05:19 -05:00
Qianfeng
9e80cdceb7 [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>
 2021-08-26 18:04:55 -07:00
zjing14
a7a758d8ce 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>
 2021-08-25 10:55:55 -05:00
zjing14
9d3f634a3c Xdlops refactor fix (#22) 
* added constexpr ahead of adptor; clean unused driver; rename M/NPerWave to M/NPerXDL

* fixed bwd

* fixed comment
 2021-08-23 11:22:10 -05:00
Chao Liu
c6f26bb480 magic division use __umulhi() (#19) 2021-08-23 10:40:27 -05:00
Chao Liu
6fe3627a9e 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>
 2021-08-19 10:55:03 -05:00
zjing14
a2ad6d3531 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
 2021-08-19 09:54:10 -05:00
zjing14
ba6f79a75e Added host_conv_wrw for verification (#15) 
* added host conv wrw
 2021-08-19 01:00:41 -05:00
Chao Liu
31b403526e Merge pull request #16 from ROCmSoftwarePlatform/develop 
Merge develop into master
 2021-08-18 11:22:34 -05:00
Chao Liu
b62bf8c3f8 Merge pull request #14 from ROCmSoftwarePlatform/miopen_downstream_init_integration 
MIOpen Downstream: Initial integration 2nd PR
 2021-08-16 16:39:40 -05:00
Chao Liu
ccc4a1d365 Merge pull request #8 from ROCmSoftwarePlatform/miopen_downstream_init_integration 2021-08-16 16:28:53 -05:00
Chao Liu
67ad47e7c1 refactor 2021-08-16 21:01:33 +00:00
Chao Liu
16effa767c refactor 2021-08-16 20:36:47 +00:00
Chao Liu
a91b68dfcb DynamicBuffer, StaticBuffer, amd_buffer_load support customized value for invalid element 2021-08-13 23:40:19 +00:00
Chao Liu
2cbabbba54 use int instead of index_t in kernel wrapper 2021-08-13 20:55:39 +00:00
Chao Liu
0834bc7635 compiler parameter use stream 2021-08-13 01:05:14 +00:00
Chao Liu
f2ac7832c6 make innner product compatiable on gfx900 2021-08-11 09:42:53 -05:00
Chao Liu
4e57b30a6a rename 2021-08-11 00:08:42 +00:00
Chao Liu
c03045ce2d rename 2021-08-10 23:45:36 +00:00
Chao Liu
b2589957f3 update CK build script 2021-08-10 22:19:13 +00:00
Chao Liu
2c48039d0e fix kernel filename 2021-08-10 22:15:23 +00:00
Chao Liu
d626dccc95 fix enum issue 2021-08-10 20:55:13 +00:00
Chao Liu
643ebd4f3e tidy 2021-08-10 07:07:11 +00:00
Chao Liu
ddd49ec9e7 fix clang warning suppression 2021-08-10 06:20:24 +00:00
Chao Liu
4f566c6221 vector/scalar pointer cast use c-style pointer cast instead of reinterpret_cast 2021-08-10 05:55:20 +00:00
Chao Liu
172036d728 add c-style pointer cast 2021-08-10 00:01:52 -05:00
Chao Liu
76f3131939 tidy 2021-08-09 18:49:59 -05:00
Chao Liu
d18428901e tidy 2021-08-09 18:20:02 -05:00
Chao Liu
f885c131d8 tidy 2021-08-09 22:13:47 +00:00
Chao Liu
80120f0a0c tidy 2021-08-09 21:10:09 +00:00
Chao Liu
c3efeb5e20 tidy 2021-08-09 19:32:07 +00:00
Chao Liu
56fc0842b3 tidy 2021-08-09 19:27:49 +00:00
Chao Liu
54fba515b3 tidy 2021-08-09 17:33:32 +00:00
Chao Liu
e62bae7a4a tidy 2021-08-09 15:11:35 +00:00
Chao Liu
24c8728942 add tidy 2021-08-08 17:41:54 +00:00
Chao Liu
61487e0a00 fix 2021-08-07 02:31:19 +00:00
Chao Liu
ae98b52ad8 remove online compilation from CK 2021-08-07 00:51:05 +00:00
Chao Liu
cb95421311 refactor 2021-08-06 22:17:51 +00:00
Chao Liu
73ca970155 Merge commit '437cc595c6e206dfebb118985b5171bbc1e29eab' into composable_kernel_init_integration_v3 2021-08-06 21:32:27 +00:00
Chao Liu
3b86646115 Merge pull request #7 from ROCmSoftwarePlatform/master 
Master
 2021-08-06 16:29:58 -05:00