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

[TIR] VNNI and ARM dot product intrinsic for tensorization #10925

Merged
merged 13 commits into from
Apr 8, 2022

Conversation

masahi
Copy link
Member

@masahi masahi commented Apr 7, 2022

Introduces a new directory python/tvm/tir/tensor_intrin where we put intrinsic descriptions written in TVMScript for various HW targets. They can be used by manual tensorized TIR schedules or auto-tensorized ones. More intrinsics, such as tensorcore ones, DP4A etc will be added later.

@junrushao1994 @vinx13 @shingjan @Hzfengsy @spectrometerHBH



@T.prim_func
def dot_product_4x4_i8i8i32_neon(
Copy link
Member Author

Choose a reason for hiding this comment

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

This is equivalent to the TE one in

def dot_int8_int8_int32_neon():
cc @tkonolige



@T.prim_func
def dot_product_4x4_i8i8i32_sdot(
Copy link
Member Author

Choose a reason for hiding this comment

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

This is equivalent to the TE one in

def dot_int8_int8_int32_neon_82(int32_lanes, dtype="uint"):



@T.prim_func
def dot_product_16x4_u8i8i32_vnni(
Copy link
Member Author

Choose a reason for hiding this comment

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

Equivalent to the TE one in

def dot_16x1x16_uint8_int8_int32_cascadelake():



# TODO(masahi): Parametrize the TVMScript description of dot product by
# shape and dtype, and share the common description with x86.
Copy link
Member Author

@masahi masahi Apr 7, 2022

Choose a reason for hiding this comment

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

cc @junrushao1994 @yelite, this is one of the common needs for meta programming support in TVMScript. I think shape parameterization is possible via specialize, but not sure if I can use that with T.Buffer syntax sugar.

A similar need arises for tensorcore (different mma shape x data type)


vec_b = B.vload([0, 0], dtype="int8x16")

# TODO(masahi): Remove duplication when inlined function call is supported
Copy link
Member Author

@masahi masahi Apr 7, 2022

Choose a reason for hiding this comment

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

cc @junrushao1994 @yelite, I want to define and call a convenience function like

def pairwise_add_mul(extract_half):
to remove duplication in IR generation. I think this is a low-hanging fruit that can be supported without more foundational work on meta-programming? (just need to tweak the python parser?)

Copy link
Member

Choose a reason for hiding this comment

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

Yes, it is likely doable and @Hzfengsy probably already has something ready

@junrushao
Copy link
Member

CC @vinx13 would you like to review? Thanks a lot!

@vinx13 vinx13 merged commit fc04738 into apache:main Apr 8, 2022
pfk-beta pushed a commit to pfk-beta/tvm that referenced this pull request Apr 11, 2022
mehrdadh pushed a commit to mehrdadh/tvm that referenced this pull request Apr 11, 2022
Lucien0 pushed a commit to Lucien0/tvm that referenced this pull request Apr 19, 2022
altanh pushed a commit to altanh/tvm that referenced this pull request Apr 28, 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.

None yet

3 participants