Dynamic Composable Kernel Initial Integration #1059
Conversation
git-subtree-dir: src/composable_kernel git-subtree-split: 1264925422920f24b3bb4fa34f178e31a23c97b5
|
@asroy > Move old static composable kernel to: src/kernels/static_composable_kernel/ Why, if we can keep old static ones as is, and add new kernels to .../dynamic_composable_kernel? |
Does it differ from |
This should not be so. We already have Int8 convolution Solvers. What is wrong? |
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 Have you looked at DeepCode's analysis? |
I plan to remove the old static kernel source code eventually |
It's used when building composable kernel alone (and outside MIOpen). It is not used in MIOpen. The whole composable kernel is pulled in as subtree |
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 |
There is not "IsInt8()" in ConvolutionContext. I tried to add it once, but compiler complain something like no "InputTensorDataType()". So I gave up |
If this is a showstopper, then please ask @DrizztDoUrden or me for assistance or open a ticket. |
…init_integration_v2
…rnel_init_integration_v2
…init_integration_v2
[Notice] If the efforts to make the composable kernels independent has resulted in the addition of another implicit dependency (a copy of addkernels), then I'm afraid that something is wrong. But well, I signed off from reviewing this, so let's actual reviewers decide. |
|
The commit history in this PR is messed up, because I did subtree pulls from different remote branches and from different composable_kernel repo forks. I realized it's important to have a single composble_kernel repo, and a single composble_kernel branch to pull from, in order to keep the commit historyu correct. I need to close this PR and open a new PR. @atamazov I will refer to all your comments here in the new PR |
[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. |
I see. Thank you. |
I don't have a good way to avoid this issue (some composable kernel host codes are not used in MIOpen) if using subtree feature. I actually brought up this issue during our 07/06 meeting as a side effect of using subtree (or submodules). The only way I can think of it's to make composable_kernel as a separately compiled library and provide a API that return a string of kernel source code to MIOpen. This was an alternative brought up by @JehandadKhan. But the in the meeting, it was decided subtree is the way to go. I'm fine with either approach, my goal is keep composable_kernel maintainable, by keeping kernel source codes in one place instead of manually copying source files between two repos. If we can reach a consensus on the best approach to achieve this goal, and if time permit, I can change accordingly |
Closing this PR. Open new one #1071 |
Add new dynamic composable kernel under:
src/composable_kernel/, which is pulled in as a "git subtree" from composable kernel repohttps://github.com/ROCmSoftwarePlatform/composable_kernel
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