Fix for distinct compile/evaluate by alexandraBara · Pull Request #533 · ROCm/MIOpen

alexandraBara · 2020-10-21T16:37:52Z

Compiling all kernels at once with Precompile kernels. This enables us to compile in parallel for Tuna.

Ran local tests to get speed for compiling/evaluating separately on MI100.
./build/bin/MIOpenDriver conv -F 2 -n 128 -g 1 -k 1024 -c 1024 -H 35 -W 35 -y 3 -x 3 -p 0 -q 0 -u 2 -v 2 -l 1 -j 1 -V 0 -w 1 -t 1 -i 1
common env vars:
export MIOPEN_FIND_ENFORCE=4
export MIOPEN_LOG_LEVEL=6
export MIOPEN_DEBUG_CONV_GEMM=0
export MIOPEN_DEBUG_CONV_FFT=0
export MIOPEN_DEBUG_CONV_DIRECT=0
export MIOPEN_DEBUG_CONV_WINOGRAD=0
export MIOPEN_DEBUG_CONV_IMPLICIT_GEMM=1

Regular compile+evaluate at the same time for igemm:
-- 5 iteration: avg 282 sec
-- 10 iteration: avg 145 sec

Separate compile, evaluate using extra env vars for compile:
export MIOPEN_COMPILE_AND_RUN=0
export MIOPEN_CUSTOM_CACHE_DIR=/root/test

First run with MIOPEN_COMPILE_PARALLEL_LEVEL=20
-- compile avg: 12 sec
-- run avg: 8 sec
-- total avg: 20 sec

Second run with MIOPEN_COMPILE_PARALLEL_LEVEL=10
-- compile avg: 19 sec
-- run avg: 8 sec
-- total avg: 27 sec

**UPDATE: parallel compile for kernels has been added as default behavior to generic_search. We can now compile in parallel and execute in the same run. No env vars needed other than set the compile level with: MIOPEN_COMPILE_PARALLEL_LEVEL=20.

We can still compile/execute separately using:
export MIOPEN_COMPILE_AND_RUN=0**

daniellowell · 2020-10-26T17:34:47Z

@alexandraBara @pfultz2 Do we understand the relationship between the multi-threading from "Find" and "Run-and-Measure"?
Can we disable parallel "Find" if we are setting tuning MIOPEN_FIND_ENFORCE > 2

zjing14 · 2020-10-27T01:11:15Z

Tuning with MIOPEN_FIND_ENFORCE=4 MIOPEN_COMPILE_AND_RUN=0 MIOPEN_CUSTOM_CACHE_DIR=/root/test MIOPEN_COMPILE_PARALLEL_LEVEL=32
Got Error Msg:
MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmForwardV4R5Xdlops: /root/MIOpen_dev/src/include/miopen/generic_search.hpp:491: Search failed
Did I use it properly?

alexandraBara · 2020-10-27T14:00:58Z

Tuning with MIOPEN_FIND_ENFORCE=4 MIOPEN_COMPILE_AND_RUN=0 MIOPEN_CUSTOM_CACHE_DIR=/root/test MIOPEN_COMPILE_PARALLEL_LEVEL=32
Got Error Msg:
MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmForwardV4R5Xdlops: /root/MIOpen_dev/src/include/miopen/generic_search.hpp:491: Search failed
Did I use it properly?

I thought the upper bound for MIOPEN_COMPILE_PARALLEL_LEVEL is 20, so i am not sure if setting it above 20 will have the effect you are thinking it would. But this wouldn't influence the outcome of the generic search.

This behavior is correct the first time around when it only compiles. We let generic_search fail so the rest of the code doesnt try to load the binaries in memory. However you should not see it fail the second time around (when you execute), after you unset MIOPEN_COMPILE_AND_RUN=0 . At this point the kernels are compiled and written to cache and the execution should continue, generic search will try and benchmark these kernels.
Maybe we can add a different error message when MIOPEN_COMPILE_AND_RUN=0 is used, something like: "Search skipped" instead of "Search failed". We can have a chat when you get online and i can explain more.

@zjing14 with the latest update to this PR you should be able to compile/run in the same execution. The only env var you need is to specify compile level threading: MIOPEN_COMPILE_PARALLEL_LEVEL=20.

…search

codecov · 2020-10-27T19:50:18Z

Codecov Report

Merging #533 into develop will increase coverage by 0.00%.
The diff coverage is 0.00%.

@@           Coverage Diff            @@
##           develop     #533   +/-   ##
========================================
  Coverage    52.22%   52.22%           
========================================
  Files          297      297           
  Lines        46020    46018    -2     
========================================
+ Hits         24032    24033    +1     
+ Misses       21988    21985    -3

Impacted Files	Coverage Δ
src/include/miopen/generic_search.hpp	`0.00% <0.00%> (ø)`
src/ocl/clhelper.cpp	`86.66% <0.00%> (-1.12%)`	⬇️
src/db_record.cpp	`89.28% <0.00%> (+1.78%)`	⬆️
src/include/miopen/find_controls.hpp	`84.00% <0.00%> (+4.00%)`	⬆️

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update d6189be...c58e17c. Read the comment docs.

daniellowell · 2020-10-28T14:50:29Z

All please rereview this PR. The update is that this has been fixed so that end users no longer need to explicitly set a flag to compiler and evaluate all at once.

atamazov · 2020-10-28T20:15:39Z

@alexandraBara

alexandraBara added the bug label 7 days ago

What bug does this PR fix?

daniellowell · 2020-10-28T22:30:46Z

@atamazov The original PR #307 made the claim that parallel compile would take place before evaluation. In fact that did not happen due to a coding bug.

atamazov

Before continuation with this PR, I highly recommend resolving review comments in the baseline: #307 (review)

atamazov · 2020-10-28T21:50:52Z

+                continue;
+            kernels.push_back(kernel);
+        }
+    }


We definitely need signs of life in this loop. Please add a new monitoring method to Heartbeat and use it here.

i added this: #if MIOPEN_ENABLE_SQLITE_KERN_CACHE around the first loop, i can add a comment

#533 (comment)

atamazov · 2020-10-28T22:31:30Z

+            kernels.push_back(kernel);
+        }
+    }
+    std::vector<Program> programs = PrecompileKernels(profile_h, kernels);


If BUILD_DEV=On, then this is waste of time, because binary cache is switched OFF.

is this a problem? we are aware that for this to work we need bin_cache turned on.

Yes, because it affects all developers that work with auto-tuning. For example, recently I tried to test auto-tuning locally and found that it takes too long, so I just stopped trying.

(I thought it was a compiler issue but actual reason is #307)

We may want to add the runtime program handles to the KernelCache (at least for development builds). That would make the loop not redundant. Alternatively, we may disable the loop when the binary cache is disabled and compile_and_run != "0".

I would prefer the simplest variant + comment about alternative solution.

i added this: #if MIOPEN_ENABLE_SQLITE_KERN_CACHE around the first loop, i can add a comment

@alexandraBara Please do this

Find IsCacheDisabled() in binary_cache.cpp

Rename it to IsBinaryCacheDisabled()

Make it global (remove static and add its declaration to binary_cache.hpp)

Use it here instead of MIOPEN_ENABLE_SQLITE_KERN_CACHE

atamazov · 2020-10-28T22:38:24Z

@daniellowell

@atamazov The original PR #307 made the claim that parallel compile would take place before evaluation. In fact that did not happen due to a coding bug.

Thanks, now I see. I thought it was intentional.

atamazov · 2020-10-28T22:39:33Z

BTW I have nothing against resolving #307 (review) right in this PR.

atamazov

I am approving this as I feel this is urgent PR. Please fix the following ASAP (possibly in some other PR):

atamazov · 2020-11-10T22:20:38Z

+#endif

-    for(const auto& current_config : all_configs)
+    if(IsEnabled(MIOPEN_DEBUG_COMPILE_ONLY{}))


🐛 Auto-tuning is not functional after this change. I am curious how this PR was tested.

* kernel parallel compile fix * added PrecompileKernels function call as default behavior in generic search

…ated to: SWDEV-304151) (#1361) * Also resolves #533 (comment)

Alex Bara added 2 commits October 16, 2020 09:33

kernel parallel compile fix

1184757

Merge branch 'develop' into alex_pcomp

e93f821

alexandraBara added the bug label Oct 21, 2020

alexandraBara requested review from DrizztDoUrden, JehandadKhan and atamazov October 21, 2020 16:37

JehandadKhan previously approved these changes Oct 21, 2020

View reviewed changes

alexandraBara added the ON_HOLD label Oct 27, 2020

Merge branch 'develop' into alex_pcomp

e0cf974

alexandraBara dismissed JehandadKhan’s stale review via 3bcd04e October 27, 2020 19:09

added PrecompileKernels function call as default behavior in generic …

3bcd04e

…search

alexandraBara removed the ON_HOLD label Oct 27, 2020

alexandraBara requested a review from JehandadKhan October 27, 2020 19:36

JehandadKhan previously approved these changes Oct 28, 2020

View reviewed changes

atamazov suggested changes Oct 28, 2020

View reviewed changes

clang tidy

b592d4f

alexandraBara dismissed JehandadKhan’s stale review via b592d4f October 29, 2020 16:56

minor changes

caf8c62

alexandraBara mentioned this pull request Oct 30, 2020

Tracking changes requested in merged PR #307 #560

Closed

Alex Bara added 2 commits October 30, 2020 14:20

cleanup

2552fcd

formatting

c58e17c

atamazov approved these changes Oct 30, 2020

View reviewed changes

daniellowell merged commit 2f3aa3d into develop Nov 1, 2020

atamazov reviewed Nov 10, 2020

View reviewed changes

alexandraBara pushed a commit that referenced this pull request Nov 16, 2020

addressed reviews from #307 and #533 (comment)

86cbc3f

JehandadKhan pushed a commit that referenced this pull request Nov 20, 2020

Fix for distinct compile/evaluate (#533)

6995384

* kernel parallel compile fix * added PrecompileKernels function call as default behavior in generic search

zjing14 pushed a commit that referenced this pull request Dec 1, 2020

Fix for distinct compile/evaluate (#533)

b8a5972

* kernel parallel compile fix * added PrecompileKernels function call as default behavior in generic search

atamazov mentioned this pull request Jan 25, 2021

Fixed invalid compilation params for fused ConvBiasActivAsm1x1U #695

Merged

atamazov deleted the alex_pcomp branch March 7, 2021 18:52

atamazov added a commit that referenced this pull request Dec 26, 2021

Fix #533 (comment)

df02daa

atamazov mentioned this pull request Dec 27, 2021

Fix for PR 1221. After tuning, remove benchmarked Programs only (related to: SWDEV-304151) #1361

Merged

atamazov added a commit that referenced this pull request Dec 29, 2021

Fix for PR #1221. After tuning, remove benchmarked Programs only (rel…

372955b

…ated to: SWDEV-304151) (#1361) * Also resolves #533 (comment)

atamazov mentioned this pull request Feb 22, 2022

Cross-Compile failing with hipErrorSharedObjectInitFailed #1439

Closed

Conversation

alexandraBara commented Oct 21, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

daniellowell commented Oct 26, 2020

Uh oh!

zjing14 commented Oct 27, 2020

Uh oh!

alexandraBara commented Oct 27, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

codecov bot commented Oct 27, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

daniellowell commented Oct 28, 2020

Uh oh!

atamazov commented Oct 28, 2020

Uh oh!

daniellowell commented Oct 28, 2020

Uh oh!

atamazov left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

DrizztDoUrden Oct 29, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

atamazov commented Oct 28, 2020

Uh oh!

atamazov commented Oct 28, 2020

Uh oh!

atamazov left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

alexandraBara commented Oct 21, 2020 •

edited

Loading

alexandraBara commented Oct 27, 2020 •

edited

Loading

codecov bot commented Oct 27, 2020 •

edited

Loading

DrizztDoUrden Oct 29, 2020 •

edited

Loading