GPU'less compile by alexandraBara · Pull Request #307 · ROCm/MIOpen

alexandraBara · 2020-06-23T16:08:36Z

I am posting this PR preemptively so Tuna people can see my changes.

The idea is:

Set env var to compile only (for desired arch/num_cu) - through Precompile kernels and dump blobs into kern_db
Move kern cache to machine where the kernels will execute
Run through regular invokers code which will compile (if kernel not already present in kern_db) and execute. Note: the kernel will already be in kern_db from the compile step.

The following steps need to happen:
build with MIOPEN_SQLITE_KERN_CACHE=ON

To compile we need to set the following env vars:
export MIOPEN_DEBUG_CONV_GEMM=0 (GEMM code path has not been implemented for this)
export MIOPEN_CUSTOM_CACHE_DIR=test (this will be appended to ~/.cache)
export MIOPEN_FIND_ENFORCE=4
export MIOPEN_COMPILE_AND_RUN=0 (compile only, escape generic search early)
export MIOPEN_DEVICE_ARCH=gfx906 (or whatever other supported arch)
export MIOPEN_DEVICE_CU=60 (or any other num_cu)

To run:
export MIOPEN_DEBUG_CONV_GEMM=0 (this can also be left out and invokers will actually compile and run as well)
export MIOPEN_CUSTOM_CACHE_DIR=test (must be same location as compile so it can find compiled kernels)
export MIOPEN_FIND_ENFORCE=4
export MIOPEN_COMPILE_AND_RUN=2 (anything but 0, can also be left unset)

I have tested as such:

compile only(for arch/num_cu different than my own machine) and check kern_db to make sure its populated with compiled kernels.
moved ukdb to machine with arch/num_cu for which it was compiled for.
set appropriate env vars to pick up custom cache location and run.
executed above steps for both hip and opencl
tested with rocm 3.5

Feel free to propose some actual utest I can add.

Blocked by: #203

to be able to override invoke type

daniellowell · 2020-09-29T15:47:02Z

 *
 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * IMPLIED, INCnLnUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,


Uhhh, what happened here?

… throw escape err

daniellowell · 2020-10-01T20:09:45Z

@DrizztDoUrden Please review this PR.

alexandraBara · 2020-10-02T14:40:30Z

@alexandraBara This change affects +989 −1,791 lines. Are all of these changes really required?

remember this is branched of ddu/invoker-tune. most of the changes are his until he merges his branch the diff will be quite big.

In this case could you please change base branch of this PR to ddu/invoker-tune? Please see https://docs.github.com/en/github/collaborating-with-issues-and-pull-requests/changing-the-base-branch-of-a-pull-request. This would greatly simplify reviewing.

I also recommend marking PR as a Draft to prevent it from merging into ddu/invoker-tune. Thanks.

Vasili's branch went in, so now its a PR from develop. Is this ok?

* Implemented igemm data invokers * Missing invoker * Removed accidental redundant unused FFT calls * Fixed kernel name split * ConvHipImplicitGemmBwdDataV4R1 invoker * Removed precompilation duplication * Removed RunAndMeasureSolution's from header * Finished merginr RaMS into invokers * Self-copy/self-move tidy fix * Fixed is_detected_v tidy warning * fixed conv_common.hpp build * Updated static_assert to use {} * env vars for custom cache location and env var to spoof arch * added PrecompileKernels fct call * calling PrecompileKernels * hip code path for compiling only, tobeadded: ocl * added env var for num_cu and removed custom cache dir * added custom bin cache location to avoid MIOPEN_VERSION_TWEAK * implemented GPU'less compiling for ocl backend * [Squashed] Generic search with invokers * Missing PrecompileSoulutions * taking bin cache location as full path, and updated gen_search to not throw escape err Co-authored-by: Vasilii Filippov <vasja.filippov94@outlook.com> Co-authored-by: Vasilii Filippov <vasilii.filippov@amd.com> Co-authored-by: daniellowell <daniel.lowell@amd.com>

atamazov · 2020-10-28T17:36:50Z

+              MIOPEN_STRINGIZE(MIOPEN_VERSION_TWEAK);
    auto p = boost::filesystem::path{miopen::ExpandUser(cache_dir)} / version;
+
+    const char* const custom = miopen::GetStringEnv(MIOPEN_CUSTOM_CACHE_DIR{});


[Quality] This is developer's feature, therefore at least DEBUG must be used in the name. And this is binary cache, let's indicate that as well.

[Recommendation] MIOPEN_DEBUG_BINARY_CACHE_DIR

atamazov · 2020-10-28T17:47:59Z

+    const char* const num_cu = miopen::GetStringEnv(MIOPEN_DEVICE_CU{});
+    if(num_cu != nullptr && strlen(num_cu) > 0)
+    {
+        return boost::lexical_cast<std::size_t>(num_cu);
+    }


[Quality]

Developers' envvar should contain DEBUG

[Recommendation] MIOPEN_DEBUG_ENFORCE_DEVICE_CU.

boost::lexical_cast should not be used.

All the checks and conversions should be moved to a function that performs the job once, to minimize delays.

[Hint] We have Value() in env.hpp for this purpose.

Partially resolved in #981

atamazov · 2020-10-28T17:56:28Z


 std::string Handle::GetDeviceName() const
 {
+    const char* const arch = miopen::GetStringEnv(MIOPEN_DEVICE_ARCH{});


[Quality] This change is not necessary

What exactly are you referring to? Returning the env var specified over the actual arch we are running on is necessary to compile for a different architecture.

We already have MIOPEN_DEBUG_ENFORCE_DEVICE. See #307 (comment)

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

+    const char* const c_and_r = miopen::GetStringEnv(MIOPEN_COMPILE_AND_RUN{});
+    std::string compile_and_run;
+    if(c_and_r != nullptr && strlen(c_and_r) > 0)
+    {
+        compile_and_run = c_and_r;
+    }
+


This is not necessary.

atamazov · 2020-10-28T21:12:06Z

        {
            current_solution = s.GetSolution(context, current_config, true);

+            if(compile_and_run == "0")


if(IsEnabled(MIOPEN_DEBUG_COMPILE_ONLY{}))

atamazov · 2020-10-28T21:18:14Z

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


[Quality] std::ignore = ...

[Note] This is just a waste of time for developers' builds (BUILD_DEV=On). Not that important, just FYI.

atamazov · 2020-10-28T21:30:25Z

+    OCLKernel(SharedProgramPtr p, const std::string& kernel_name)
+        : program(p), kernel(CreateKernel(p.get(), kernel_name))
+    {
+    }
+


[Quality] Do not add a new ctor. Just update existing one: read MIOPEN_DEBUG_COMPILE_ONLY and skip loading and running the kernels, if enabled.

atamazov · 2020-10-28T21:34:54Z

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


[Quality] You may wish adding a new monitoring method to Heartbeat and using it here, to show signs of life when log level < 6.

* Implemented igemm data invokers * Missing invoker * Removed accidental redundant unused FFT calls * Fixed kernel name split * ConvHipImplicitGemmBwdDataV4R1 invoker * Removed precompilation duplication * Removed RunAndMeasureSolution's from header * Finished merginr RaMS into invokers * Self-copy/self-move tidy fix * Fixed is_detected_v tidy warning * fixed conv_common.hpp build * Updated static_assert to use {} * env vars for custom cache location and env var to spoof arch * added PrecompileKernels fct call * calling PrecompileKernels * hip code path for compiling only, tobeadded: ocl * added env var for num_cu and removed custom cache dir * added custom bin cache location to avoid MIOPEN_VERSION_TWEAK * implemented GPU'less compiling for ocl backend * [Squashed] Generic search with invokers * Missing PrecompileSoulutions * taking bin cache location as full path, and updated gen_search to not throw escape err Co-authored-by: Vasilii Filippov <vasja.filippov94@outlook.com> Co-authored-by: Vasilii Filippov <vasilii.filippov@amd.com> Co-authored-by: daniellowell <daniel.lowell@amd.com>

* Resolve #307 (comment) and #307 (comment) * Partially resolve #307 (comment) * Resolve Resolve #307 (comment) * Remove W/A for https://ontrack-internal.amd.com/browse/SWDEV-225285

Vasilii Filippov and others added 30 commits April 9, 2020 14:56

Implemented igemm data invokers

a76133f

format

322bcf5

Tidy fix

4613f4b

Reverted scgemm function header to pre-merge state

8c639bb

Missing invoker

2f20f7f

Merge remote-tracking branch 'public/develop' into ddu/igemm-invokers

0229117

Removed accidental redundant unused FFT calls

307b081

Merge remote-tracking branch 'origin/develop' into ddu/igemm-invokers

6d54f45

Merge remote-tracking branch 'origin/develop' into ddu/igemm-invokers

50f11c0

Fixed kernel name split

d5afb6d

EOL spaces

da096cb

Fixed build after merge

acec640

clang-format

6067335

EoL spaces

7839970

Includes order

66679cc

ConvHipImplicitGemmBwdDataV4R1 invoker

dcef7f8

fixed the build

cc7268e

Removed precompilation duplication

c2eb69a

WIP

d44ae26

WIP: invoke_ctx passing everywhere for auto-tune

6f39360

Merge remote-tracking branch 'origin/develop' into ddu/invoker-tune

ad03566

WIP: Updated generic search to use invokers

18f17d7

Removed RunAndMeasureSolution's from header

33cfbc8

WIP: boost::any -> AnyInvokeParams

7481e50

to be able to override invoke type

WIP: Merged RunAndMeasure into most invokers

f1fc387

Cleaned-up implicit gemm

eafafff

Finished merginr RaMS into invokers

72e534c

Comments and cleanup

ce2f61b

Merge remote-tracking branch 'origin/develop' into ddu/invoker-tune

fdbc37a

format

ed7ae93

daniellowell reviewed Sep 29, 2020

View reviewed changes

taking bin cache location as full path, and updated gen_search to not…

5ced2f6

… throw escape err

Alex Bara added 2 commits October 2, 2020 09:43

cleanup

65e5bd7

Merge branch 'develop' into alex_parallel_compile

dbf5c91

JehandadKhan requested changes Oct 5, 2020

View reviewed changes

Comment thread src/binary_cache.cpp

Comment thread src/hipoc/hipoc_kernel.cpp

Comment thread src/ocl_kernel.cpp

Alex Bara added 2 commits October 5, 2020 11:19

clang format fix

44a2817

more format fix

b158703

JehandadKhan previously approved these changes Oct 5, 2020

View reviewed changes

more clang fix

fc1ef9c

alexandraBara dismissed JehandadKhan’s stale review via fc1ef9c October 5, 2020 20:40

alexandraBara changed the title ~~Preemptive PR on GPU'less compile~~ GPU'less compile Oct 6, 2020

daniellowell merged commit 6fd907b into develop Oct 7, 2020

atamazov reviewed Oct 28, 2020

View reviewed changes

daniellowell mentioned this pull request Oct 28, 2020

Fix for distinct compile/evaluate #533

Merged

alexandraBara mentioned this pull request Oct 30, 2020

Tracking changes requested in merged PR #307 #560

Closed

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

addressed reviews from #307 and #533 (comment)

86cbc3f

JehandadKhan mentioned this pull request Jan 20, 2021

No GPU backend for Tuna #681

Merged

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

atamazov added a commit that referenced this pull request Jun 11, 2021

Resolve #307 (comment) and #307 (comment)

02c80de

atamazov added a commit that referenced this pull request Jun 11, 2021

Partially resolve #307 (comment)

0a1d50a

atamazov added a commit that referenced this pull request Jun 11, 2021

Resolve Resolve #307 (comment)

3e2bce0

atamazov mentioned this pull request Jun 11, 2021

Remove W/A for SWDEV-225285 #981

Merged

atamazov mentioned this pull request Feb 22, 2022

Cross-Compile failing with hipErrorSharedObjectInitFailed #1439

Closed

Conversation

alexandraBara commented Jun 23, 2020 • edited by JehandadKhan Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

daniellowell commented Oct 1, 2020

Uh oh!

alexandraBara commented Oct 2, 2020

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

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!

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!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

alexandraBara commented Jun 23, 2020 •

edited by JehandadKhan

Loading