-
Notifications
You must be signed in to change notification settings - Fork 3.5k
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
Conversation
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 |
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 |
There was a problem hiding this comment.
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?
There was a problem hiding this comment.
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.
There was a problem hiding this comment.
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
There was a problem hiding this comment.
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.
This PR adds support for the dynamic sequence length in cutlass attention.
cc @cyx-6 @masahi