MIOpenTensile Part1 : Add MIOpenTensile in GEMM path #139

ce1adon · 2020-04-08T00:07:23Z

The purpose of this PR is to add MIOpenTensile path in GEMM algorithm for fp32 on hip backend. It replaces the default GEMM path through rocblas in HIP backend for fp32 type.

ghost · 2020-04-08T00:07:38Z

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

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

atamazov · 2020-04-08T01:48:54Z

The Invokers work is complicated by frequent changes in the baseline code. For example, recent additions of ImplicitGemm solvers have created substantial additional work for @DrizztDoUrden, and, thus have delayed the overall Invoker's ETA.

@DrizztDoUrden is also reworking Gemm code to Solvers and Invokers. This work is currently paused because Vasya is doing something more important (ImplicitGemm Invokers), but I heard it is almost ready. If this PR will be merged first, then its code will need to be reworked once again almost immediately after merging, and that is we all would like to avoid, I guess.

I am not aware of the priorities of this MIOpenTensile thing. @daniellowell @ce1adon @DrizztDoUrden could you please communicate and decide the optimal order?

DrizztDoUrden · 2020-04-09T13:07:44Z

Doesn't matter much unless it's conv...ocl.cpp changes, which are not here. I don't change gemm internals there.

atamazov · 2020-04-09T13:15:16Z

Okay, then. @ce1adon Could you please set priority label.

atamazov · 2020-04-29T13:58:57Z

@ce1adon Please describe how the PR affects the library's behavior. Or this is NFC?

ce1adon · 2020-04-29T18:46:36Z

@ce1adon Please describe how the PR affects the library's behavior. Or this is NFC?

It's not a new feature. The influence is updated in top comment.

ce1adon · 2020-05-19T16:42:49Z

Jira ticket related to the blocker of this PR
http://ontrack-internal.amd.com/browse/SWDEV-233998

atamazov · 2020-09-25T23:10:22Z

@ce1adon Up to you. However, AFAIU on_hold does not prevent anything. It simply says that PR should not be merged.

…upport bf16/fp16

src/solver/conv_multipass_wino3x3WrW.cpp

pfultz2 · 2020-09-28T23:02:47Z

hip-clang.docker

+RUN export HIPCC_LINK_FLAGS_APPEND='-O3 -parallel-jobs=4'
+RUN export HIPCC_COMPILE_FLAGS_APPEND='-O3 -Wno-format-nonliteral -parallel-jobs=4'
+# install last released miopentensile in default, install latest commits when MIOTENSILE_VER="latest"
+RUN if [ "$MIOTENSILE_VER" = "latest" ] ; then cget -p $PREFIX install -DAMDGPU_TARGETS=${GPU_ARCH} ROCmSoftwarePlatform/MIOpenTensile@e258bbaed0bd9de546ea38c9a5c42f71fa41d9a0; else cget -p $PREFIX install -DAMDGPU_TARGETS=${GPU_ARCH} ROCmSoftwarePlatform/MIOpenTensile@2e3e792b2674bf8cdf2620749a298ed5313351bb; fi


-DAMDGPU_TARGETS=${GPU_ARCH} should be passed cget init. Also this should install tensile from a file so that hcc and hip-clang install the same version.

Will move gpu_arch to init.
Currently both hcc and hip-clang are using the same version. The version divergence here is both for hip-clang. It intends to test both a stable version and a latest version of yaml+tensile.

moved gpu_arch flag to cget init

…-clang dockerfile

ce1adon · 2020-10-09T17:31:36Z

The last test passed but is marked as failed.

ce1adon · 2020-10-26T17:51:25Z

NOTE: d0969a0 passed at # 38

atamazov

LGTM

* add miopentensile path * use miotensile in hip backend * update miotensile * generalize miotensile stride * disable miopentensile in ocl backend * remove ocl backend consecutive identical branches * disable use of miotensile in hip-clang * update miopentensile version * use const arg in gemm * miot introduce transpose feature * update miot version * disable miotensile in hip-clang in cmakelists * disable miotensile except in hcc in cmakelists * disable ocl wrw2 solvers * enable miopentensile built with hipclang * add miotensile in hip-clang docker * add python3-venv in hip-clang docker * hip-clang docker update * update hip-clang docker compiler * update miopentensile info * disable miotensile with hcc * set miopentensile off & add dedicated test stage * Revert "clang tidy" This reverts commit bd706c5. * Revert "disable ocl wrw2 solvers" This reverts commit 17f5a59. * workaround for issue 2534 * Adjust test config * Adjust gemm env_var * use miotensile in winograd3x3multipass solver * revert gemm enforce backend values * skip igemm tests in test stage iv (miotensile tests) * skip igemm tests in test stage iv * use parallel jobs for tensile * specify gpu architecture(s) in miotensile build process * fix jenkins docker arg * fix jenkins docker arg string * rename gpu target param for miotensile * specify gpu architecture(s) in Jenkinsfile buildJob * relocate miotensile gpu targets arg * add symlink to /opt/rocm in Dockerfile * revise Dockerfile to pass GPU architecture flag for dependencies * add jenkins tests for last/latest version of miopentensile * remove workaround for issue #2534 * winograd workaround condition update: current miopentensile doesn't support bf16/fp16 * fix jenkinsfile typo & move GPU architecture flag to cget init in hip-clang dockerfile * revise conv_multipass_wino3x3WrW GEMM workaround for miopentensile * miopentensile stage rename * reset miopentensile version in hcc * revert wino3x3WrW datatype constraint in miopentensile * remove miopentensile from dependency list & update version Co-authored-by: Jing Zhou <zhou@ixt-rack-95.local.lan>

atamazov · 2021-02-18T18:59:09Z

test/CMakeLists.txt

    OR (MIOPEN_TEST_BFLOAT16 AND ${PARSE_ALLOW_BFLOAT16})
    OR (MIOPEN_TEST_HALF AND ${PARSE_ALLOW_HALF})
-    OR (MIOPEN_TEST_INT8 AND ${PARSE_ALLOW_INT8}))
+    OR (MIOPEN_TEST_INT8 AND ${PARSE_ALLOW_INT8})
+    OR (NOT (MIOPEN_TEST_MIOTENSILE AND (NAME IN_LIST SKIP_TESTS))))


@ce1adon It should be

OR (MIOPEN_TEST_MIOTENSILE AND NOT (NAME IN_LIST SKIP_TESTS)))

Current code enables ALL tests when MIOPEN_TEST_MIOTENSILE=off, thus affecting non-miopentensile tests (half etc). I will fix that.

This is under the condition of (NOT (MIOPEN_TEST_INT8 OR MIOPEN_TEST_BFLOAT16 OR MIOPEN_TEST_HALF OR MIOPEN_TEST_MIOTENSILE)), which means this only applies to fp32 cases. Do you mean fp32 doesn't need to test all the cases?

No, this is not "under" that confition. It is just one of ORed conditions. Alternate formatting:

if( (NOT (MIOPEN_TEST_INT8 OR MIOPEN_TEST_BFLOAT16 OR MIOPEN_TEST_HALF OR MIOPEN_TEST_MIOTENSILE)) // (1) OR (MIOPEN_TEST_BFLOAT16 AND ${PARSE_ALLOW_BFLOAT16}) // (2) OR ( MIOPEN_TEST_HALF AND ${PARSE_ALLOW_HALF}) // (3) OR (MIOPEN_TEST_INT8 AND ${PARSE_ALLOW_INT8}) // (4) OR (NOT (MIOPEN_TEST_MIOTENSILE AND (NAME IN_LIST SKIP_TESTS))) // (5) ) ...

When MIOPEN_TEST_MIOTENSILE=Off, condition (5) is always true. This works properly for MIOpenTensile, that's why you've missed that, I guess.

~~Fixed in 318f480 (not yet merged to develop)~~

Fixed in #750

atamazov · 2021-11-06T12:25:08Z

@ce1adon

It replaces the default GEMM path through rocblas in HIP backend for fp32 type.

If MIOpenTensile is the default for FP32, then shall we remove FP32 testing from "MIOpenTensile" stages in the Jenkinsfile?

add miopentensile path

3201498

ce1adon requested review from pfultz2, atamazov and daniellowell April 8, 2020 00:07

atamazov requested a review from DrizztDoUrden April 8, 2020 01:19

use miotensile in hip backend

46daceb

ce1adon force-pushed the miotensile branch from d86a2cc to 46daceb Compare April 9, 2020 00:29

ce1adon added the value_high label Apr 9, 2020

hip tidy

bb28471

ce1adon force-pushed the miotensile branch from c16414d to 8ca947f Compare April 10, 2020 22:22

update miotensile

3c6d5bd

ce1adon force-pushed the miotensile branch from 8ca947f to 3c6d5bd Compare April 11, 2020 02:16

ce1adon added 6 commits April 13, 2020 17:06

generalize miotensile stride

a0d8cbc

disable miopentensile in ocl backend

93e51f2

remove ocl backend consecutive identical branches

0aa1a64

disable use of miotensile in hip-clang

e00580c

Merge branch 'develop' into miotensile

583dd50

update miopentensile version

736940c

Merge branch 'develop' into miotensile

5e29d30

Merge branch 'develop' into miotensile

33b0fbd

atamazov added the ON_HOLD label May 25, 2020

ce1adon added 2 commits June 21, 2020 16:14

Merge branch 'develop' into miotensile

0998746

use const arg in gemm

818544f

ce1adon removed the ON_HOLD label Sep 25, 2020

Jing Zhou added 4 commits September 26, 2020 00:49

add jenkins tests for last/latest version of miopentensile

67a06c7

remove workaround for issue #2534

569bd7b

winograd workaround condition update: current miopentensile doesn't s…

6b28286

…upport bf16/fp16

clang format

8d50d0e

ce1adon dismissed atamazov’s stale review via 8d50d0e September 26, 2020 01:22

Merge branch 'develop' into miotensile

20b6e44

atamazov suggested changes Sep 28, 2020

View reviewed changes

src/solver/conv_multipass_wino3x3WrW.cpp Show resolved Hide resolved

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

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

pfultz2 reviewed Sep 28, 2020

View reviewed changes

ce1adon and others added 7 commits September 29, 2020 01:25

fix jenkinsfile typo & move GPU architecture flag to cget init in hip…

29c0524

…-clang dockerfile

revise conv_multipass_wino3x3WrW GEMM workaround for miopentensile

402b02e

clang format

99581d7

miopentensile stage rename

0e2dd07

reset miopentensile version in hcc

b49fee4

revert wino3x3WrW datatype constraint in miopentensile

4d7cc64

Merge branch 'develop' into miotensile

ee6b154

ce1adon added 2 commits October 10, 2020 03:30

remove miopentensile from dependency list & update version

3402052

Merge branch 'develop' into miotensile

d0969a0

atamazov approved these changes Oct 26, 2020

View reviewed changes

ce1adon changed the title ~~Add MIOpenTensile in GEMM path~~ MIOpenTensile Part1 : Add MIOpenTensile in GEMM path Oct 29, 2020

ce1adon mentioned this pull request Oct 29, 2020

MIOpenTensile Part2 : Support low precision types #556

Merged

daniellowell merged commit 6e91fe5 into develop Oct 30, 2020

ce1adon deleted the miotensile branch November 14, 2020 23:02

atamazov reviewed Feb 18, 2021

View reviewed changes

atamazov mentioned this pull request Feb 23, 2021

[tests] Remove useless tests from low-precision stages. #750

Merged

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

MIOpenTensile Part1 : Add MIOpenTensile in GEMM path #139

MIOpenTensile Part1 : Add MIOpenTensile in GEMM path #139

ce1adon commented Apr 8, 2020 •

edited

Loading

ghost commented Apr 8, 2020 •

edited by ghost

Loading

atamazov commented Apr 8, 2020

DrizztDoUrden commented Apr 9, 2020

atamazov commented Apr 9, 2020

atamazov commented Apr 29, 2020

ce1adon commented Apr 29, 2020

ce1adon commented May 19, 2020 •

edited

Loading

atamazov commented Sep 25, 2020

pfultz2 Sep 28, 2020

ce1adon Sep 28, 2020

ce1adon Sep 29, 2020

ce1adon commented Oct 9, 2020

ce1adon commented Oct 26, 2020

atamazov left a comment

atamazov Feb 18, 2021

ce1adon Feb 18, 2021

atamazov Feb 18, 2021 •

edited

Loading

atamazov Feb 18, 2021 •

edited

Loading

atamazov Feb 22, 2021

atamazov commented Nov 6, 2021

MIOpenTensile Part1 : Add MIOpenTensile in GEMM path #139

MIOpenTensile Part1 : Add MIOpenTensile in GEMM path #139

Conversation

ce1adon commented Apr 8, 2020 • edited Loading

ghost commented Apr 8, 2020 • edited by ghost Loading

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

atamazov commented Apr 8, 2020

DrizztDoUrden commented Apr 9, 2020

atamazov commented Apr 9, 2020

atamazov commented Apr 29, 2020

ce1adon commented Apr 29, 2020

ce1adon commented May 19, 2020 • edited Loading

atamazov commented Sep 25, 2020

pfultz2 Sep 28, 2020

Choose a reason for hiding this comment

ce1adon Sep 28, 2020

Choose a reason for hiding this comment

ce1adon Sep 29, 2020

Choose a reason for hiding this comment

ce1adon commented Oct 9, 2020

ce1adon commented Oct 26, 2020

atamazov left a comment

Choose a reason for hiding this comment

atamazov Feb 18, 2021

Choose a reason for hiding this comment

ce1adon Feb 18, 2021

Choose a reason for hiding this comment

atamazov Feb 18, 2021 • edited Loading

Choose a reason for hiding this comment

atamazov Feb 18, 2021 • edited Loading

Choose a reason for hiding this comment

atamazov Feb 22, 2021

Choose a reason for hiding this comment

atamazov commented Nov 6, 2021

ce1adon commented Apr 8, 2020 •

edited

Loading

ghost commented Apr 8, 2020 •

edited by ghost

Loading

ce1adon commented May 19, 2020 •

edited

Loading

atamazov Feb 18, 2021 •

edited

Loading

atamazov Feb 18, 2021 •

edited

Loading