Skip to content

[TIR] Preserve annotations after lower opaque block#12572

Merged
Hzfengsy merged 2 commits intoapache:mainfrom
wrongtest-intellif:support_preserve_customize_annotations
Aug 31, 2022
Merged

[TIR] Preserve annotations after lower opaque block#12572
Hzfengsy merged 2 commits intoapache:mainfrom
wrongtest-intellif:support_preserve_customize_annotations

Conversation

@wrongtest-intellif
Copy link
Contributor

@wrongtest-intellif wrongtest-intellif commented Aug 24, 2022

Hi, there. The change aims to improve customization experience of TIR lowering using annotations.

Currently, we have LowerOpaqueBlock pass to convert s-tir(schedulable tir with blocks) to non-stir. All annotations are dropped then except those used by legacy TE pragma attributes, they are converted to legacy AttrStmt. For example:

# before
for i in range(10, annotations={"pragma_myattr0": 1, "my_attr0": 1}): ....

# after
with T.attr(i, "pragma_myattr0", 1):
    for i in range(10): ....

This introduce difficulties when we want to leverage schedule phase annotations in lowering passes.

  1. pragma_ annotations are converted to legacy AttrStmt, which only accepts PrimExpr typed value. But the loop/block annotations before are originally much flexible, one could use list/dict to encode more rich compile time hints.
  2. Using AttrStmt loses the roundtrip property of TIR, thus we fail to write testcases of lowering pass with friendly before/after script comparations.
  3. It is more cohesive to encode annotations just with related IR node (like ForNode). Or else the developers always have to write separate logics handling attr node and for node.

The change just add a pass configuration "tir.LowerOpaqueBlock": {"preserved_annotations": List[str]}}. For each annotation key value of loop or block:

  1. if key starts with pragma_, keep the original behaviour.
  2. if key is registered in preserved_annotations list, the annotation is preserved.
  3. or else it is dropped.

The alternative may be preserve all annotations without any extra configuration. But I think current way take the least impact on overall lowering stream. If we do not add customized pass config into pass context, all things should be same.

cc @Hzfengsy @junrushao1994

@github-actions github-actions bot requested a review from Hzfengsy August 24, 2022 11:15
@T.prim_func
def after(A: T.Buffer[8, "float32"], B: T.Buffer[8, "float32"]):
for i in T.serial(8, annotations={"k_0": 1, "k_1": [2, 3]}):
for _ in range(1, annotations={"k_3": 3.14}):
Copy link
Member

Choose a reason for hiding this comment

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

It‘s a bit tricky to add a unit loop here. Can we keep the current behavior or add an AttrStmt?

Copy link
Member

Choose a reason for hiding this comment

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

I'm a bit uncertain if haivng an AttrStmt might make sense in our workflow, given it's something in the low-level TIR that we might want to minimize its usage (CC: @tqchen). I'm aware that there could be any implications if we directly remove a unit loop without preserving its annotation too. Not sure what the best approach is.

Copy link
Member

@tqchen tqchen Aug 24, 2022

Choose a reason for hiding this comment

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

On AttrStmt, we should move away from it when possible as an effort to limit the scope and clarify the semantics. Of course, we should also document the set of possible annotatons we use in For so they become manageable


mod = tvm.IRModule.from_expr(before)
with tvm.transform.PassContext(
config={"tir.LowerOpaqueBlock": {"preserved_annotations": ["k_0", "k_1", "k_3"]}}
Copy link
Member

Choose a reason for hiding this comment

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

Is it possible to preserve all annotations by default rather than write a pass config?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

A tracing of usages in current test_tir_ and test_meta_schedule_ show that for loop annotations, only software pipeline specified keys are used, they are consumed before lower opaque block. The other keys are mostly meta_schedule.* in block annotations.

What about that for non-pragma keys?

  • Always preserve loop annotations.

  • Always drop block annotations, since they are mostly designed for meta-schedule thus make no sense to subsequent passes, and it seems not clear now how to lower the block's attrs after blocks get eliminated.

For my circumstance the loop annotations are quite useful while block annotations are not used for lowering passes. Could the decision about block annotation get delayed to when it is indeed needed :)?

Copy link
Member

Choose a reason for hiding this comment

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

@Hzfengsy Hzfengsy merged commit f114d55 into apache:main Aug 31, 2022
xinetzone pushed a commit to daobook/tvm that referenced this pull request Nov 25, 2022
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