Dynamic Composable Kernel Initial Integration #1071

asroy · 2021-07-30T23:02:20Z

PR description

Add new dynamic composable kernel under:
src/composable_kernel/, which is pulled in as a "git subtree" from composable kernel repo
remote repo: https://github.com/ROCmSoftwarePlatform/composable_kernel
remote branch: master

Move old static composable kernel to:
src/kernels/static_composable_kernel/

Add a new iGEMM solver for Forward NCHW FP32/FP16/Int8 (int8 is not supported by ConvolutionContext so the solver support for Int8 cannot be used yet):
src/solver/conv_ck_igemm_fwd_v6r1_dlops_nchw.cpp::ConvCkIgemmFwdV6r1DlopsNchw , which is calls into composable kernel solver:
src/composable_kernel/host/driver_online/include/conv_igemm_fwd_v6r1_dlops_nchw_kcyx_nkhw.hpp::ConvIgemmFwdV6r1DlopsNchwKcyxNkhw

Sync between `composable_kernel` and `MIOpen` with "git subtree"

Relation between MIOpen solver class `ConvCkIgemmFwdV6r1DlopsNchw` and composable_kernel class `ConvIgemmFwdV6r1DlopsNchwKcyxNkhw`

…sable_kernel'

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

… files

ghost · 2021-07-30T23:02:27Z

Congratulations 🎉. DeepCode analyzed your code in 2.66 seconds and we found no issues. Enjoy a moment of no bugs ☀️.

👉 View analysis in DeepCode’s Dashboard | Configure the bot

👉 The DeepCode service and API will be deprecated in August, 2021. Here is the information how to migrate. Thank you for using DeepCode 🙏 ❤️ !

If you are using our plugins, you might be interested in their successors: Snyk's JetBrains plugin and Snyk's VS Code plugin.

asroy · 2021-07-30T23:10:52Z

Previous review comment

[@atamazov] #1059 (comment)

Add new dynamic composable kernel under: src/composable_kernel/, which is pulled in as a "git subtree" from composable kernel repo

I am sorry but I do not remember the motivation for subtree... Where can I refresh my memory (mail thread or ticket in the repo...)?

[@asroy] #1059 (comment)

The motivation is to be able to sync composable kernel repo and MIOpen automatically.

We had the meeting in 07/06 "MIOpen and composable kernel", between you, @JehandadKhan @pfultz2, me and @zjing14. Subtree approach was suggested there

[@atamazov] #1059 (comment)

[Notice] Thanks for the explanation. Maybe I do not understand the meaning of "automatically" in this context. But would it allow something like a one-liner PR in MIOpen to completely change the implementation of the Solvers that employ composable kernels? If yes, then this looks dangerous to me.

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: 5781adf5cf4ac753e2e36da7385791775b744bf7

…le_kernel_init_integration_v3

asroy · 2021-08-07T02:45:41Z

I've removed online compilation code from this PR.
@atamazov @JehandadKhan @junliume

asroy · 2021-08-11T22:52:00Z

CI has passed. This PR is ready to be reviewed

junliume · 2021-08-11T23:06:23Z

CI has passed. This PR is ready to be reviewed

Ping @pfultz2 @zjing14 @JehandadKhan @qianfengz for review :)

src/include/miopen/solver.hpp

src/include/miopen/solver/ck_utility_common.hpp

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

… for invalid element

…init_integration_v3

src/solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops_padded_gemm.cpp

src/CMakeLists.txt

src/composable_kernel/composable_kernel/include/utility/static_buffer.hpp

src/include/miopen/solver/ck_utility_common.hpp

qianfengz · 2021-08-18T13:10:20Z

src/composable_kernel/composable_kernel/include/utility/dynamic_buffer.hpp

+
+template <AddressSpaceEnum_t BufferAddressSpace, typename T, typename ElementSpaceSize>
+__host__ __device__ constexpr auto
+make_dynamic_buffer(T* p, ElementSpaceSize element_space_size, T invalid_element_value)


Should we use remove_cv_t<T> here to provide a simpler type for invalid_element_value, so that the calling codes don't have to make difficult adaption

I will update this in CK repo master branch. Reduction PR will be able to use the update

asroy · 2021-08-18T22:52:33Z

@junliume @JehandadKhan CI just passed

junliume · 2021-08-18T23:29:58Z

@junliume @JehandadKhan CI just passed

@asroy would you need to make more code changes? e.g. comments from #1071 (comment)

If not, any objections that we merge it tonight? (2000 CDT, 1800 PDT)

asroy · 2021-08-19T14:11:39Z

@junliume No more change is needed for this PR.

Chao Liu added 3 commits July 30, 2021 17:31

Merge commit 'c840438b62e3071b8e658de7343c8e461387de97' as 'src/compo…

78c293c

…sable_kernel'

Squashed 'src/composable_kernel/' content from commit f6edda611

c840438

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

add solver ConvIgemmFwdV6r1DlopsNchwKcyxNkhw; rename static ck source…

6204be8

… files

asroy added value_high IMPACTS_PERFORMANCE_DB IMPACTS_FIND_DB complexity_high new feature initial iteration time urgency_high labels Jul 30, 2021

asroy requested review from pfultz2, zjing14, JehandadKhan, qianfengz and atamazov July 30, 2021 23:02

This comment has been minimized.

Sign in to view

asroy mentioned this pull request Jul 30, 2021

Dynamic Composable Kernel Initial Integration #1059

Closed

Chao Liu added 6 commits August 6, 2021 21:32

Merge commit '437cc595c6e206dfebb118985b5171bbc1e29eab' into composab…

3c0614b

…le_kernel_init_integration_v3

fix

d822708

refactor

0bf90d4

remove online compilation from CK

c67b040

refactor

fff94ec

asroy mentioned this pull request Aug 7, 2021

MIOpen Downstream: Initial integration ROCm/composable_kernel#8

Merged

fix

44da477

add ctest

1208667

make innner product compatiable on gfx900

4619a4e

junliume added TESTING_CI_PASSED urgency_blocker labels Aug 11, 2021

JehandadKhan requested changes Aug 12, 2021

View reviewed changes

src/include/miopen/solver.hpp Show resolved Hide resolved

src/include/miopen/solver/ck_utility_common.hpp Outdated Show resolved Hide resolved

src/include/miopen/solver/ck_utility_common.hpp Outdated Show resolved Hide resolved

Chao Liu and others added 3 commits August 12, 2021 17:39

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

d5e7530

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

compiler parameter use stream

04999e3

use int instead of index_t in kernel wrapper

b460246

asroy requested a review from JehandadKhan August 13, 2021 20:56

Chao Liu added 5 commits August 13, 2021 23:40

DynamicBuffer, StaticBuffer, amd_buffer_load support customized value…

6445cae

… for invalid element

Merge remote-tracking branch 'origin/develop' into composable_kernel_…

40860ed

…init_integration_v3

Merge remote-tracking branch 'origin/develop' into composable_kernel_…

09f375e

…init_integration_v3

refactor

18a7570

refactor

3241ab4

qianfengz reviewed Aug 17, 2021

View reviewed changes

Chao Liu added 2 commits August 17, 2021 10:33

change cmakelist

d50dc04

change ck common utility

8db34a5

JehandadKhan previously approved these changes Aug 17, 2021

View reviewed changes

fix

77b4c43

asroy dismissed JehandadKhan’s stale review via 77b4c43 August 17, 2021 23:51

qianfengz reviewed Aug 18, 2021

View reviewed changes

junliume approved these changes Aug 18, 2021

View reviewed changes

JehandadKhan merged commit da5a391 into develop Aug 19, 2021

JehandadKhan mentioned this pull request Aug 19, 2021

Revert "Dynamic Composable Kernel Initial Integration" #1096

Closed

asroy mentioned this pull request Aug 19, 2021

Copy of "Dynamic Composable Kernel Initial Integration #1071" #1097

Merged

JehandadKhan deleted the composable_kernel_init_integration_v3 branch August 19, 2021 15:59

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Dynamic Composable Kernel Initial Integration #1071

Dynamic Composable Kernel Initial Integration #1071

asroy commented Jul 30, 2021 •

edited

Loading

ghost commented Jul 30, 2021 •

edited by ghost

Loading

asroy commented Jul 30, 2021 •

edited

Loading

This comment has been minimized.

This comment has been minimized.

asroy commented Aug 7, 2021 •

edited

Loading

asroy commented Aug 11, 2021

junliume commented Aug 11, 2021

qianfengz Aug 18, 2021

asroy Aug 18, 2021

asroy commented Aug 18, 2021

junliume commented Aug 18, 2021

asroy commented Aug 19, 2021

Dynamic Composable Kernel Initial Integration #1071

Dynamic Composable Kernel Initial Integration #1071

Conversation

asroy commented Jul 30, 2021 • edited Loading

PR description

Sync between composable_kernel and MIOpen with "git subtree"

Relation between MIOpen solver class ConvCkIgemmFwdV6r1DlopsNchw and composable_kernel class ConvIgemmFwdV6r1DlopsNchwKcyxNkhw

ghost commented Jul 30, 2021 • edited by ghost Loading

👉 View analysis in DeepCode’s Dashboard | Configure the bot

👉 The DeepCode service and API will be deprecated in August, 2021. Here is the information how to migrate. Thank you for using DeepCode 🙏 ❤️ !

asroy commented Jul 30, 2021 • edited Loading

This comment has been minimized.

This comment has been minimized.

asroy commented Aug 7, 2021 • edited Loading

asroy commented Aug 11, 2021

junliume commented Aug 11, 2021

qianfengz Aug 18, 2021

Choose a reason for hiding this comment

asroy Aug 18, 2021

Choose a reason for hiding this comment

asroy commented Aug 18, 2021

junliume commented Aug 18, 2021

asroy commented Aug 19, 2021

asroy commented Jul 30, 2021 •

edited

Loading

Sync between `composable_kernel` and `MIOpen` with "git subtree"

Relation between MIOpen solver class `ConvCkIgemmFwdV6r1DlopsNchw` and composable_kernel class `ConvIgemmFwdV6r1DlopsNchwKcyxNkhw`

ghost commented Jul 30, 2021 •

edited by ghost

Loading

asroy commented Jul 30, 2021 •

edited

Loading

asroy commented Aug 7, 2021 •

edited

Loading