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
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 f6edda611f6edda61195781adf5c#5 ) (#6 )
97e6d514f#4 from ROCmSoftwarePlatform/separate_online_compile
7b1ec41e549c33aaea54b3e73d15781adf5cfJehandad.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
Chao Liu
5781adf5cf
Update develop ( #5 ) ( #6 )
...
* refactor
2021-08-06 16:13:07 -05:00
Chao Liu
d09ea4f4e5
Update develop ( #5 )
...
* refactor
2021-08-06 16:11:15 -05:00
