You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by "masahi (via GitHub)" <gi...@apache.org> on 2023/05/07 20:13:55 UTC
[GitHub] [tvm] masahi opened a new pull request, #14798: [CUTLASS] Add NDEBUG option to CUTLASS compile to speed up attention kernel
masahi opened a new pull request, #14798:
URL: https://github.com/apache/tvm/pull/14798
I found that adding this option brings non-trivial perf improvement to the attention kernel (2.4 vs 2.7 msec for the most heavy workload in SD UNet). This results in a few msec speed up for SD UNet e2e.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] vinx13 commented on pull request #14798: [CUTLASS] Add NDEBUG option to CUTLASS compile to speed up attention kernel
Posted by "vinx13 (via GitHub)" <gi...@apache.org>.
vinx13 commented on PR #14798:
URL: https://github.com/apache/tvm/pull/14798#issuecomment-1540980726
triton kernel is also causal attention, this mask https://github.com/openai/triton/blob/main/python/tutorials/06-fused-attention.py#LL55C25-L55C25 will result in less computation
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] vinx13 commented on pull request #14798: [CUTLASS] Add NDEBUG option to CUTLASS compile to speed up attention kernel
Posted by "vinx13 (via GitHub)" <gi...@apache.org>.
vinx13 commented on PR #14798:
URL: https://github.com/apache/tvm/pull/14798#issuecomment-1540999409
triton kernel is also causal attention, https://github.com/openai/triton/blob/main/python/tutorials/06-fused-attention.py#LL49C57-L49C57 it's not doing a full Q*K.T, so it has less computation
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #14798: [CUTLASS] Add NDEBUG option to CUTLASS compile to speed up attention kernel
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on PR #14798:
URL: https://github.com/apache/tvm/pull/14798#issuecomment-1538075959
I tried updating the cutlass submodule revision, but it didn't help.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tvm-bot commented on pull request #14798: [CUTLASS] Add NDEBUG option to CUTLASS compile to speed up attention kernel
Posted by "tvm-bot (via GitHub)" <gi...@apache.org>.
tvm-bot commented on PR #14798:
URL: https://github.com/apache/tvm/pull/14798#issuecomment-1537531328
<!---bot-comment-->
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](https://github.com/apache/incubator-tvm/blob/master/CONTRIBUTORS.md#reviewers) by @-ing them in a comment.
<!--bot-comment-ccs-start-->
* No users to tag found in teams: `cutlass` <sub>See [#10317](https://github.com/apache/tvm/issues/10317) for details</sub><!--bot-comment-ccs-end-->
<sub>Generated by [tvm-bot](https://github.com/apache/tvm/blob/main/ci/README.md#github-actions)</sub>
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] vinx13 merged pull request #14798: [CUTLASS] Add NDEBUG option to CUTLASS compile to speed up attention kernel
Posted by "vinx13 (via GitHub)" <gi...@apache.org>.
vinx13 merged PR #14798:
URL: https://github.com/apache/tvm/pull/14798
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #14798: [CUTLASS] Add NDEBUG option to CUTLASS compile to speed up attention kernel
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on PR #14798:
URL: https://github.com/apache/tvm/pull/14798#issuecomment-1537543264
But a weird thing is, if I run the same attention workload via the cutlass example, it shows that the same kernel runs in 1.3 msec, see below (compared to our BYOC result, 2.4 msec).
```
$ nsys nvprof examples/41_fused_multi_head_attention/41_fused_multi_head_attention_fixed_seqlen --head_number=8 --batch_size=2 --head_size=40 --_head_size_v=40 --seq_length=4096 --seq_length_kv=4096
CUTLASS Attention: ====================================================
{seq length Q, seq length KV, head size, head size V, head number, batch size} = {4096, 4096, 40, 40, 8, 2}.
Runtime: 1.36964 ms
GFLOPs: 19897.7
Passed
Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) GridXYZ BlockXYZ
Name
-------- --------------- --------- ----------- ----------- --------- --------- ----------- -------------- -------------- ------------------
...
13.8 30,160,277 22 1,370,921.7 1,368,225.0 1,346,113 1,414,209 16,718.9 64 8 2 32 4 1 void attention_kernel_batched_impl<AttentionKernel<cutlass::half_t, cutlass::arch::Sm80, (bool)1, (…
```
I've also checked out Triton and Flash attention perf on the same workload by running https://github.com/openai/triton/blob/main/python/tutorials/06-fused-attention.py#L330 and saw perf around 1.2 - 1.3 msec. So I want to believe that ~1.3 msec should be the right result for an attention kernel on this workload.
So maybe there is something off in how we use this kernel from our BYOC? I compared [the generated code](https://gist.github.com/masahi/767e84ee17e9648621ff48251fa5bba5) and the cutlass example code but didn't find any difference. There is not difference in the nvcc options that might affect performance other than this `NDEBUG` stuff (that's how I found about it). Any thoughts? @vinx13 @cyx-6
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #14798: [CUTLASS] Add NDEBUG option to CUTLASS compile to speed up attention kernel
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on PR #14798:
URL: https://github.com/apache/tvm/pull/14798#issuecomment-1540978051
wow then Triton and flash attention kernels may indeed be faster than the cutlass one, given that the triton implementation is definitely not doing causal mask optimization.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] junrushao commented on pull request #14798: [CUTLASS] Add NDEBUG option to CUTLASS compile to speed up attention kernel
Posted by "junrushao (via GitHub)" <gi...@apache.org>.
junrushao commented on PR #14798:
URL: https://github.com/apache/tvm/pull/14798#issuecomment-1537545283
@spectrometerHBH has same observation. Good catch!
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] vinx13 commented on pull request #14798: [CUTLASS] Add NDEBUG option to CUTLASS compile to speed up attention kernel
Posted by "vinx13 (via GitHub)" <gi...@apache.org>.
vinx13 commented on PR #14798:
URL: https://github.com/apache/tvm/pull/14798#issuecomment-1538007758
I can confirm there's performance difference, the profiler also shows different number of instructions are executed, though the it's indeed the same kernel
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] vinx13 commented on pull request #14798: [CUTLASS] Add NDEBUG option to CUTLASS compile to speed up attention kernel
Posted by "vinx13 (via GitHub)" <gi...@apache.org>.
vinx13 commented on PR #14798:
URL: https://github.com/apache/tvm/pull/14798#issuecomment-1540028748
it turns out that cutlass profiler has `--causal` default to true, after adding `--causal=false` I can get the same numbers
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] masahi commented on pull request #14798: [CUTLASS] Add NDEBUG option to CUTLASS compile to speed up attention kernel
Posted by "masahi (via GitHub)" <gi...@apache.org>.
masahi commented on PR #14798:
URL: https://github.com/apache/tvm/pull/14798#issuecomment-1541008894
Interesting! Yeah I didn't understand this loop bound `for start_n in range(0, (start_m + 1) * BLOCK_M, BLOCK_N)` when I was studying this code, now it makes sense.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org