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

Disable DMA transpose on attention forward kernel #42

Open
wants to merge 1 commit into
base: main
Choose a base branch
from

Conversation

aws-zhehongb
Copy link

Issue #, if available:

Description of changes:

Disable DMA transpose on attention forward kernel while we are stabilizing the DMA transpose in trn2

Testing:

Please see detailed unit test requirements in the CONTRIBUTING.md

  • [ *] The change is covered by numeric check using nki.baremetal
  • [ *] The change is covered by performance benchmark test using nki.benchmark
  • [* ] The change is covered by end-to-end integration test

Pull Request Checklist

  • [ *] I have filled in all the required field in the template
  • [* ] I have tested locally that all the tests pass
  • [* ] By submitting this pull request, I confirm that my contribution is made under the terms of the MIT-0 license.

While we are stabilizing the DMA transpose
Copy link

github-actions bot commented Jan 9, 2025

Please download and review the generated API doc at https://github.com/aws-neuron/nki-samples/actions/runs/12696429583

@@ -44,20 +44,13 @@ class FlashConfig:
@nki.jit(mode='trace')
def transpose_p_local(p_local_transposed, p_local, LARGE_TILE_SZ):
for i in nl.affine_range(LARGE_TILE_SZ // 512):
if nisa.get_nc_version() == nisa.nc_version.gen3:
Copy link
Contributor

Choose a reason for hiding this comment

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

could we have a comment on why is the DMA transpose is disabled?

@aws-qieqingy
Copy link
Contributor

Internal testing for the PR passed. Ready to merge with another approval.

@aws-qieqingy aws-qieqingy requested a review from ggumen January 20, 2025 17:28
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.

2 participants