Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Finish support for list-of-targets #11382

Merged
merged 12 commits into from
May 23, 2022

Conversation

mbs-octoml
Copy link
Contributor

@mbs-octoml mbs-octoml commented May 19, 2022

This finishes the work started in #11173 to support
'external codegen' targets in the N build-like API surfaces, which in turn requires all the
Relay-level build interfaces to support list-of-targets instead of dictionary-of-targets. And
prepare for plumbing targets through the external codegen machinery.

(See https://discuss.tvm.apache.org/t/byoc-supporting-cutlass-byoc-with-collage/12796/6 for
more context on build/target changes, and https://github.com/apache/tvm-rfcs/blob/main/rfcs/0062-collage.md
for the overall Collage RFC).

  • It turns out it's ok if a build is given only a single 'external codegen' target, so remove that check
    in CompilationConfig::Init. (When Collage builds a 'candidate partition' it does so for a single target.
    As far as Collage is concerned it does not care whether the target is regular (eg Target("cuda")), or
    for a specific external codegen (eg Target("cutlass")), it just passes the target into the build.)

  • Add CompilationConfig::FindPrimitiveTargetForKind which I'll later need to retrieve
    the external codegen Target instance corresponding to a "Compiler" attribute value.

  • Target.update_target_host_consist was supporting three API styles:

    • single target
    • map from device type to target
    • map from target to IRModule (for the ir_to_runtime API)

    I replaced all those calls with a more specialized 'canonicalize' call:

    • Target.canonicalize_target_and_host
    • Target.canonicalize_multi_targets_and_host
    • Target.canonicalize_target_map_and_host

    In particular, all the tuning interfaces (task extraction, tuning, tuning records) all explicitly
    do not support multiple targets since that just doesn't make sense.

  • Rev the micro library format version from 5 to 6 since the 'target' attribute in the metadata.json
    output is now an array rather than a dictionary. Remove the N uses of target.values() to follow
    suite.

assert len(raw_targets) == 1
assert raw_targets[0].kind.name == "cuda"
assert raw_targets[0].host.kind.name == "llvm"


def test_canon_multi_target_and_host_6():
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Have you thought about structuring tests into classes according to the function under test?
Like
class TestCanonMultiTargetAndHost:
def test_with_tvm_objects(self):
# this test here

The names of tests can then be shorter and more descriptive at the same time.
This is more a question than a suggestion.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Not really experienced with good pytest style, just mimicking whatever I happen to find already in place. Happy to try that if it's considered the preferred idiom?

This finishes the work started in apache#11173 to support
'external codegen' targets in the N build-like API surfaces.

 - It turns out it's ok if a build is given only a single 'external codegen' target, so remove that check
   in CompilationConfig::Init. When Collage builds a 'candidate partition' it does so for a single target.
   As far as Collage is concerned it does not care whether the target is regular (eg Target("cuda")), or
   for a specific external codegen (eg Target("cutlass")), it just passes the target into the build.

 - Add CompilationConfig::FindPrimitiveTargetForKind which I'll later need to retrieve
   the external codegen Target instance corresponding to a "Compiler" attribute value.

 - Target.update_target_host_consist was supporting three API styles:
    - single target
    - map from device type to target
    - map from target to IRModule (for the ir_to_runtime API)
   I replaced all those calls with a more specialized 'canonicalize' call:
    - Target.canonicalize_target_and_host
    - Target.canonicalize_multi_targets_and_host
    - Target.canonicalize_target_map_and_host
   In particular, all the tuning interfaces (task extraction, tuning, tuning records) all explicitly
   *do not* support multiple targets since the underlying code just doesn't support that.
- Revert unintended changes
- Improve comments in compilation_config.h
- Update target/target_host params documentation
- Rev micro library format from 5 to 6
- Use Target.current() in a few places
- Handle host already being in Target
- Take device type from target
Copy link
Contributor

@jwfromm jwfromm left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This overall LGTM. Thank you @mbs-octoml!

@jwfromm jwfromm merged commit d146777 into apache:main May 23, 2022
@mbs-octoml mbs-octoml deleted the mbs-collage-more-targets branch May 23, 2022 16:52
@vinx13
Copy link
Member

vinx13 commented May 24, 2022

Hi @mbs-octoml with this PR I'm seeing some warnings UserWarning: target_host parameter is going to be deprecated. Please pass in tvm.target.Target(target, host=target_host) instead. even if I don't use target_host. Would you mind taking a look? It can be reproduced with the following script

import tvm
from tvm import te
target = tvm.target.Target('nvidia/geforce-rtx-3070')
N = 32
A = te.placeholder(name='A', shape=[N])
B = te.compute((N,), lambda i,: A[i] + 1.0)
sch = te.create_schedule(B.op)
sch[B].bind(sch[B].op.axis[0], te.thread_axis('threadIdx.x'))
tvm.build(sch, [A, B], target)

@mbs-octoml
Copy link
Contributor Author

mbs-octoml commented May 25, 2022

#11499

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

4 participants