Flash Attention for Neuron #883
Open
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
apivovarov
reviewed
Dec 11, 2024
apoorvtintin
force-pushed
the
mainline_upstream_fa
branch
from
eab90f9 to
347f522
Compare
kelvin-zou
reviewed
Dec 11, 2024
| @@ -0,0 +1,129 @@ | |||
| from absl import logging | |||
There was a problem hiding this comment.
Also please add comments for the file.
| return out | ||
|
|
||
|
|
||
| def _mha_forward(query, key, value, bias, causal, softmax_scale): |
There was a problem hiding this comment.
Can we get a support for segment ID and dropout as well? Both are quite needed nowadays.
| @@ -159,6 +159,21 @@ def jit_attn(query, key, value, bias, segment_ids): | |||
|
|
|||
| return jit_attn | |||
|
|
|||
| elif backend == "neuron": | |||
| from axlearn.common.flash_attention.neuron_attention import ( | |||
There was a problem hiding this comment.
On demand import is kind of risky, we can live with it for functions inside neuron_attention.py, can we at least get it as a header import for files outside of neuron_attention.py?
| k = key.transpose(0, 2, 3, 1) # [batch_size, num_heads, d_model, kv_seq_len] | ||
| v = value.transpose(0, 2, 1, 3) # [batch_size, num_heads, kv_seq_len, d_model] | ||
|
|
||
| import neuronxcc.nki.language as nl |
ruomingp
reviewed
Dec 14, 2024
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This PR adds support for flash attention kernel for Neuron implemented through Neuron Kernel Interface (NKI).
The flash attention kernel works with TRN1 and TRN2.