[MLIR] Binary backend - step 1 - nonxdlops fwd path #841

jerryyin · 2021-04-02T14:38:41Z

This PR enables the back-to-back execution of mlir binary backend fwd kernels.

I did the following in this PR:

Bumped the MLIR commit to latest
Implemented the MLIR binary backend invoker
Implemented the MLIR binary backend solver
- Added .mlir extension for binary backend compilation
- Amended the call chain such that both .mlir-cpp and .mlir extension are handled
Pulled in memref.hpp header from a subset of llvm source, refer to here
Added unit test for mlir binary backend

Please note: CI in full pass with initial commit 5ea65fb.

Also this is a follow-up from @atamazov 's review in this commit.

Update: MLIR repository has a breaking change in group convolution support. Because of that, I wouldn't be able to address any MLIR end concerns without adding the new feature. Therefore I will work in the following sequence:

Get the first round of fwd/bwd/wrw solvers implemented and submitted to MIOpen
Submit Adding StridedMemRefType from interface rocMLIR#136
Bring MLIR to the latest, support group convolution, drop StridedMemRef

src/mlir_build.cpp

atamazov · 2021-04-09T15:15:24Z

Reviewers, please review ;)

src/include/miopen/memref.hpp

test/CMakeLists.txt

shurale-nkn

LGTM

atamazov · 2021-04-09T22:45:36Z

[Off-topic] @jerryyin Please try to not forget to terminate outdated CI job after pushing to the repo. Thanks!

jerryyin · 2021-04-12T14:25:36Z

@atamazov Sounds good, will keep a closer eye on delinquent CI jobs.

atamazov

It seems like we should not borrow StridedMemRefType<> from MLIR. Conceptually, it represents the layout of MLIR kernel parameters. IIUC this layout may change in the future, so the relevant abstraction should be a part of the Miir API. Otherwise we will have a situation when some part of Miir implementation is penetrated into MIOpen.

Luckily we do not need much. The only thing we need is a struct type (POD type equivalent to the data members of StridedMemRefType<void, 4>). You can define it in the Miir.h, thus isolating MIOpen from the specific layout of MLIR kernel parameters.

When/if the implementation of StridedMemRefType<> in MLIR change, you would need to update that struct type in Miir.h.

Sounds reasonable?

src/conv/invokers/mlir_impl_gemm.cpp

atamazov · 2021-04-12T21:38:13Z

src/conv/invokers/mlir_impl_gemm.cpp

+        target.basePtr = const_cast<void*>(ptr); // NOLINT (cppcoreguidelines-pro-type-const-cast)
+        target.data    = const_cast<void*>(ptr); // NOLINT (cppcoreguidelines-pro-type-const-cast)


[Optimization] Only initialization of pointers (buffer addresses) should be left in the Invoker... (see next comment)

Let's postpone this. #862 created.

atamazov · 2021-04-12T21:44:43Z

src/conv/invokers/mlir_impl_gemm.cpp

+        std::copy(dims.cbegin(), dims.cend(), &target.sizes[0]);
+        std::copy(strides.cbegin(), strides.cend(), &target.strides[0]);


...while strides and sizes and known in the Factory (outer lambda), and necessary initializations should be performed there. In this case the only thing that Invoker will need to do (before launching the kernel) is initialization of buffer addresses.

@jerryyin Whatever design you are going to select, it should be able to satisfy this and this. This is the simplest test of design validity.

Let's postpone this. #862 created.

src/include/miopen/memref.hpp

atamazov · 2021-04-18T20:36:56Z

@JehandadKhan #795 merged in. Please make necessary changes right in this PR to resolve #795 (comment). If you do not have time, please ask me and I will do that.

@jerryyin Plase let us know if any objections.

atamazov · 2021-04-18T20:37:35Z

@jerryyin Merge conflicts. Please let me know if you need help.

JehandadKhan · 2021-04-19T15:05:06Z

@JehandadKhan #795 merged in. Please make necessary changes right in this PR to resolve #795 (comment). If you do not have time, please ask me and I will do that.

@atamazov I have updated the branch.

* Fix find_permutation * Misc fixes

atamazov · 2021-04-19T21:50:43Z

One step to the approval left: #841 (comment)

atamazov

LGTM!

src/solver/conv_hip_implicit_gemm_mlir_bin_fwd.cpp

atamazov

LGTM

Implemented the fwd invoker for mlir bin backend kernels

5ea65fb

jerryyin requested review from pfultz2, DrizztDoUrden and atamazov April 2, 2021 14:38

This comment has been minimized.

Sign in to view

jerryyin commented Apr 2, 2021

View reviewed changes

src/mlir_build.cpp Outdated Show resolved Hide resolved

jerryyin added TESTING_CI_PASSED urgency_high labels Apr 6, 2021

Merge branch 'develop' into igemm-bin2

7513b8d

jerryyin force-pushed the igemm-bin2 branch 4 times, most recently from f0316ea to 153a8f4 Compare April 8, 2021 20:30

Refactor invoker factory

6128704

jerryyin force-pushed the igemm-bin2 branch from 153a8f4 to 6128704 Compare April 8, 2021 20:56

jerryyin requested review from JehandadKhan and shurale-nkn April 9, 2021 14:59

atamazov added the value_high label Apr 9, 2021

shurale-nkn requested changes Apr 9, 2021

View reviewed changes

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

test/CMakeLists.txt Show resolved Hide resolved

jerryyin requested a review from shurale-nkn April 9, 2021 17:08

shurale-nkn requested changes Apr 9, 2021

View reviewed changes

test/CMakeLists.txt Show resolved Hide resolved

Shrink the test coverage to MLIR test only in its stage

a69e45a

shurale-nkn previously approved these changes Apr 9, 2021

View reviewed changes

jerryyin changed the title ~~Implemented the fwd invoker and solver for mlir binary backend~~ Implement mlir binary backend - step 1: fwd path Apr 12, 2021

atamazov suggested changes Apr 12, 2021

View reviewed changes

jerryyin commented Apr 13, 2021

View reviewed changes

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

Move non pointer initialization to outer lambda

066cf18

atamazov mentioned this pull request Apr 17, 2021

Additions to the HIPNOGPU backend #795

Merged

atamazov removed the TESTING_CI_PASSED label Apr 18, 2021

jerryyin and others added 2 commits April 19, 2021 09:53

Merge branch 'develop' into igemm-bin2

d12724b

update load program guard for no-gpu backend

ca5da3a

Addressing review feedbacks

8fc3457

* Fix find_permutation * Misc fixes

atamazov mentioned this pull request Apr 19, 2021

[Recommendation] ConvHipImplicitGemmMlirBinFwd -> ConvMlirImplicitGemmBinFwd and ditto all other MLIR solvers. #869

Closed

Removing FAST_HYBRID specific flag from solver

113fab9

Removing CalculateGemmSize() from public solver interface

8821782

atamazov previously approved these changes Apr 19, 2021

View reviewed changes

Suppressing tidy failures

d3424e0

jerryyin dismissed atamazov’s stale review via d3424e0 April 20, 2021 00:03

atamazov reviewed Apr 20, 2021

View reviewed changes

src/solver/conv_hip_implicit_gemm_mlir_bin_fwd.cpp Outdated Show resolved Hide resolved

jerryyin added 2 commits April 20, 2021 00:50

Remove suppresion, add MIOPEN_USE_MLIR

57a7189

Merge remote-tracking branch 'origin/develop' into igemm-bin2

52dedf7

jerryyin mentioned this pull request Apr 22, 2021

[MLIR] Deprecate MLIR CPP backend - step 1 #879

Merged

jerryyin added the TESTING_CI_PASSED label Apr 23, 2021

atamazov approved these changes Apr 23, 2021

View reviewed changes

atamazov merged commit b9b5337 into develop Apr 23, 2021

This was referenced Apr 23, 2021

[MLIR] Binary backend - step 2 - nonxdlops bwd path #889

Merged

[MLIR] Binary backend - step 3 - noxdlops WrW path #892

Merged

jerryyin changed the title ~~Implement mlir binary backend - step 1: fwd path~~ [MLIR] Binary backend - step 1 - fwd path May 13, 2021

jerryyin changed the title ~~[MLIR] Binary backend - step 1 - fwd path~~ [MLIR] Binary backend - step 1 - nonxdlops fwd path May 13, 2021

This was referenced May 13, 2021

[MLIR] Binary backend - step 6 - Introduce GetKernelCount API #927

Merged

[MLIR] Binary backend - step 7 - Implement v4r1 algorithm for bwd solvers #935

Merged

jerryyin mentioned this pull request Jun 9, 2021

[MLIR] Binary backend - step 8/of 8 - Bring up and enable bwd solvers #978

Merged

atamazov deleted the igemm-bin2 branch June 23, 2021 13:19

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[MLIR] Binary backend - step 1 - nonxdlops fwd path #841

[MLIR] Binary backend - step 1 - nonxdlops fwd path #841

jerryyin commented Apr 2, 2021 •

edited

Loading

This comment has been minimized.

atamazov commented Apr 9, 2021

shurale-nkn left a comment

atamazov commented Apr 9, 2021

jerryyin commented Apr 12, 2021

atamazov left a comment •

edited

Loading

atamazov Apr 12, 2021

atamazov Apr 17, 2021

atamazov Apr 12, 2021

atamazov Apr 15, 2021

atamazov Apr 17, 2021

atamazov commented Apr 18, 2021

atamazov commented Apr 18, 2021

JehandadKhan commented Apr 19, 2021

atamazov commented Apr 19, 2021

atamazov left a comment

atamazov left a comment

		target.basePtr = const_cast<void*>(ptr); // NOLINT (cppcoreguidelines-pro-type-const-cast)
		target.data = const_cast<void*>(ptr); // NOLINT (cppcoreguidelines-pro-type-const-cast)

		std::copy(dims.cbegin(), dims.cend(), &target.sizes[0]);
		std::copy(strides.cbegin(), strides.cend(), &target.strides[0]);

[MLIR] Binary backend - step 1 - nonxdlops fwd path #841

[MLIR] Binary backend - step 1 - nonxdlops fwd path #841

Conversation

jerryyin commented Apr 2, 2021 • edited Loading

This comment has been minimized.

atamazov commented Apr 9, 2021

shurale-nkn left a comment

Choose a reason for hiding this comment

atamazov commented Apr 9, 2021

jerryyin commented Apr 12, 2021

atamazov left a comment • edited Loading

Choose a reason for hiding this comment

atamazov Apr 12, 2021

Choose a reason for hiding this comment

atamazov Apr 17, 2021

Choose a reason for hiding this comment

atamazov Apr 12, 2021

Choose a reason for hiding this comment

atamazov Apr 15, 2021

Choose a reason for hiding this comment

atamazov Apr 17, 2021

Choose a reason for hiding this comment

atamazov commented Apr 18, 2021

atamazov commented Apr 18, 2021

JehandadKhan commented Apr 19, 2021

atamazov commented Apr 19, 2021

atamazov left a comment

Choose a reason for hiding this comment

atamazov left a comment

Choose a reason for hiding this comment

jerryyin commented Apr 2, 2021 •

edited

Loading

atamazov left a comment •

edited

Loading