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

[Unity] Cutlass attention with dynamic sequence length #15028

Merged
merged 2 commits into from
Jun 6, 2023

Conversation

yelite
Copy link
Contributor

@yelite yelite commented Jun 5, 2023

This PR adds support for the dynamic sequence length in cutlass attention.

cc @cyx-6 @masahi

@tvm-bot
Copy link
Collaborator

tvm-bot commented Jun 5, 2023

Thanks for contributing to TVM! Please refer to the contributing guidelines https://tvm.apache.org/docs/contribute/ for useful information and tips. Please request code reviews from Reviewers by @-ing them in a comment.

Generated by tvm-bot

@github-actions github-actions bot requested review from cyx-6 and masahi June 5, 2023 20:03
if not isinstance(workspace_size_bytes, (int, tvm.tir.expr.IntImm)):
# Tempororay workaround for dynamic shape workload. Will be removed when
# workspace for dynamic shape workload is implemented.
workspace_size_bytes = 8
Copy link
Member

Choose a reason for hiding this comment

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

How hard would it to be to support workspace alloc for dynamic attention?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

The approach in my mind is to make the max_workspace_size_ of WorkspaceProvider a PrimExpr isntead of IntImm, which consists of a series of max expr to get the actual max workspace size from symbolic shape variables.

I tried to doing this in the most straightforward way and got This IR is not well formed: Symbolic Var a presents in different functions in the same Module.. Apparently it needs more sophisticated analysis to replace the shape variables with valid ones as the workspace size is propagated from composite function to its caller. Also I am not sure if the arg of alloc_tensor can be a non-constant expr too.

Copy link
Contributor

Choose a reason for hiding this comment

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

Just for curiosity, what is the rationale behind workspace_size_bytes = 8?

The approach in my mind is to make the max_workspace_size_ of WorkspaceProvider a PrimExpr isntead of IntImm, which consists of a series of max expr to get the actual max workspace size from symbolic shape variables.

Would you elaborate a little more? Did you try to do more sophisticated approach than using the annotation like R.func_attr({"tir_var_upper_bound": {"m": 2048, "n": 2048}})?

I am not sure if the arg of alloc_tensor can be a non-constant expr too.

aloc_tensor can take symbolic shape. See https://github.com/apache/tvm/blob/unity/tests/python/relax/test_transform_static_plan_block_memory.py#L958

Copy link
Member

@masahi masahi Jun 6, 2023

Choose a reason for hiding this comment

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

workspace_size_bytes = 8

It is just for a temp workaround, it will cause cudaMalloc to be always called.

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