Skip to content

Faster split of QKV for FlashAttention - #166

Merged
ptrendx merged 10 commits into
NVIDIA:mainfrom
ptrendx:pr_faster_qkvsplit
Apr 26, 2023
Merged

Faster split of QKV for FlashAttention#166
ptrendx merged 10 commits into
NVIDIA:mainfrom
ptrendx:pr_faster_qkvsplit

Conversation

@ptrendx

@ptrendxptrendx commented Apr 22, 2023

Copy link
Copy Markdown
Member

No description provided.

@ptrendx
Faster split of QKV for FlashAttention
dc28e6c
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx

ptrendx commented Apr 22, 2023

Copy link
Copy Markdown
MemberAuthor

/te-ci

@ptrendx
CI fixes
0d00228
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx

ptrendx commented Apr 22, 2023

Copy link
Copy Markdown
MemberAuthor

/te-ci

@ptrendx
Fix
cb0569a
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx

ptrendx commented Apr 22, 2023

Copy link
Copy Markdown
MemberAuthor

/te-ci

ksivaman
ksivaman reviewed Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman reviewed Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman reviewed Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman reviewed Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman reviewed Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Comment on lines +1910 to +1934
if (q.scalar_type() == at::ScalarType::Half) {
using dtype = at::Half;
flash_attention::prepare_kernel_bwd<dtype><<<grid, threads, 0,
at::cuda::getCurrentCUDAStream()>>>(
q.data_ptr<dtype>(),
k.data_ptr<dtype>(),
v.data_ptr<dtype>(),
qkv.data_ptr<dtype>(),
q.size(0),
q.size(1),
q.size(2),
q.size(3));
} else {
using dtype = at::BFloat16;
flash_attention::prepare_kernel_bwd<dtype><<<grid, threads, 0,
at::cuda::getCurrentCUDAStream()>>>(
q.data_ptr<dtype>(),
k.data_ptr<dtype>(),
v.data_ptr<dtype>(),
qkv.data_ptr<dtype>(),
q.size(0),
q.size(1),
q.size(2),
q.size(3));
}

@ksivamanksivamanApr 22, 2023

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

https://github.com/pytorch/pytorch/blob/main/aten/src/ATen/Dispatch.h/#L284 seems like a cleaner method

@ptrendxptrendxApr 23, 2023

Copy link
Copy Markdown
MemberAuthor

Choose a reason for hiding this comment

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

Will look into it in a follow up generalization PR (since in principle this code should work for any type, not just FP16 and BF16).

ksivaman
ksivaman reviewed Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu Outdated
ksivaman
ksivaman reviewed Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py Outdated
ksivaman
ksivaman reviewed Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
ksivamanand others added 3 commits April 24, 2023 10:02
@ksivaman
Merge branch 'main' into pr_faster_qkvsplit
b32b1ee
@ksivaman
review comments
249c370
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@ksivaman@ptrendx
Message with assert
0a6510e
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10 reviewed Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu Outdated
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
@ksivaman
Review comments
41ba126
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10 approved these changes Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py Outdated
@ksivaman
review
7508bee
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@ksivaman
fix misalignment error
e04160c
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10 requested changes Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
@ksivaman
make clarifying comment and check strides
41b515f
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
ksivaman
ksivaman approved these changes Apr 24, 2023
View reviewed changes

@ksivamanksivaman left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

LGTM

@ksivaman

ksivaman commented Apr 24, 2023

Copy link
Copy Markdown
Member

/te-ci

@ptrendx
ptrendx merged commit 8adb1b1 into NVIDIA:mainApr 26, 2023
ksivaman added a commit that referenced this pull request Apr 28, 2023
@ksivaman
Faster split of QKV for FlashAttention (#166)
71488db
* Faster split of QKV for FlashAttention
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* CI fixes
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* Fix
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Message with assert
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix misalignment error
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* make clarifying comment and check strides
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
vasunvidia pushed a commit to vasunvidia/TransformerEngine that referenced this pull request May 12, 2023
@ptrendx@ksivaman@vasunvidia
Faster split of QKV for FlashAttention (NVIDIA#166)
2ffedae
* Faster split of QKV for FlashAttention
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* CI fixes
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* Fix
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Message with assert
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix misalignment error
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* make clarifying comment and check strides
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@hungryGeek16hungryGeek16 mentioned this pull request May 31, 2026
Open
Sign up for freeto join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@timmoon10timmoon10timmoon10 approved these changes
@ksivamanksivamanksivaman approved these changes

Assignees

No one assigned

Labels

0.8.0

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

3 participants

@ptrendx@ksivaman@timmoon10
, 'i'); if (__m === '*' || __re.test(location.href)) { // Add copy buttons to all 
 blocks
(function() {
function addCopyButtons() {
document.querySelectorAll('pre code').forEach(function(codeBlock) {
if (codeBlock.parentElement.hasAttribute('data-copy-added')) return;
codeBlock.parentElement.setAttribute('data-copy-added', 'true');
var btn = document.createElement('button');
btn.textContent = 'Copy';
btn.style.cssText = 'position:absolute;top:4px;right:4px;padding:2px 8px;font-size:11px;background:#4ecdc4;border:none;border-radius:4px;color:#1a1a2e;cursor:pointer;opacity:0.7;transition:opacity 0.2s;';
btn.onmouseover = function() { this.style.opacity = '1'; };
btn.onmouseout = function() { this.style.opacity = '0.7'; };
btn.onclick = function() {
navigator.clipboard.writeText(codeBlock.textContent).then(function() {
btn.textContent = 'Copied!';
setTimeout(function() { btn.textContent = 'Copy'; }, 1500);
});
};
codeBlock.parentElement.style.position = 'relative';
codeBlock.parentElement.appendChild(btn);
});
}
addCopyButtons();
// Re-run on dynamic content
var observer = new MutationObserver(addCopyButtons);
observer.observe(document.body, { childList: true, subtree: true });
})();
}
} catch(__e) { console.warn('[Userscript:Add Copy Buttons to Code Blocks]', __e); }
})();
(function(){
try {
var __m = "github.com";
var __re = new RegExp('^' + "github\\.com" + '
Faster split of QKV for FlashAttention by ptrendx · Pull Request #166 · NVIDIA/TransformerEngine · GitHub
Skip to content
Faster split of QKV for FlashAttention - #166
Merged
ptrendx merged 10 commits into
NVIDIA:mainfrom 
ptrendx:pr_faster_qkvsplit
Apr 26, 2023
Merged
Faster split of QKV for FlashAttention#166
ptrendx merged 10 commits into
NVIDIA:mainfrom 
ptrendx:pr_faster_qkvsplit
Conversation
@ptrendx
@ptrendxptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

Member
 No description provided. 
@ptrendx
Faster split of QKV for FlashAttention
dc28e6c
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx
ptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

MemberAuthor
/te-ci
@ptrendx
CI fixes
0d00228
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx
ptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

MemberAuthor
/te-ci
@ptrendx
Fix
cb0569a
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx
ptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

MemberAuthor
/te-ci
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu

Comment on lines

+1910
 to 
+1934

 if (q.scalar_type() == at::ScalarType::Half) {
 using dtype = at::Half;
 flash_attention::prepare_kernel_bwd<dtype><<<grid, threads, 0,
 at::cuda::getCurrentCUDAStream()>>>(
 q.data_ptr<dtype>(),
 k.data_ptr<dtype>(),
 v.data_ptr<dtype>(),
 qkv.data_ptr<dtype>(),
 q.size(0),
 q.size(1),
 q.size(2),
 q.size(3));
 } else {
 using dtype = at::BFloat16;
 flash_attention::prepare_kernel_bwd<dtype><<<grid, threads, 0,
 at::cuda::getCurrentCUDAStream()>>>(
 q.data_ptr<dtype>(),
 k.data_ptr<dtype>(),
 v.data_ptr<dtype>(),
 qkv.data_ptr<dtype>(),
 q.size(0),
 q.size(1),
 q.size(2),
 q.size(3));
 }
@ksivamanksivamanApr 22, 2023

Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

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

https://github.com/pytorch/pytorch/blob/main/aten/src/ATen/Dispatch.h/#L284 seems like a cleaner method
@ptrendxptrendxApr 23, 2023

Copy link


Copy Markdown

MemberAuthor
Choose a reason for hiding this comment

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

Will look into it in a follow up generalization PR (since in principle this code should work for any type, not just FP16 and BF16).
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
Outdated

ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
ksivamanand others
added 3 commits
April 24, 2023 10:02
@ksivaman
Merge branch 'main' into pr_faster_qkvsplit
b32b1ee
@ksivaman
review comments
249c370
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@ksivaman@ptrendx
Message with assert
0a6510e
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10
reviewed
Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.cu
@ksivaman
Review comments
41ba126
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10
approved these changes
Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
Outdated

@ksivaman
review
7508bee
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@ksivaman
fix misalignment error
e04160c
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10
requested changes
Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
@ksivaman
make clarifying comment and check strides
41b515f
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
ksivaman
ksivaman
approved these changes
Apr 24, 2023
View reviewed changes
@ksivamanksivaman
left a comment


Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

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

LGTM
@ksivaman
ksivaman
commented
Apr 24, 2023

Copy link


Copy Markdown

Member
/te-ci
@ptrendx
ptrendx
merged commit 8adb1b1
into
NVIDIA:mainApr 26, 2023
ksivaman added a commit
that referenced
this pull request
Apr 28, 2023
@ksivaman
Faster split of QKV for FlashAttention (#166)
71488db
* Faster split of QKV for FlashAttention
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* CI fixes
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* Fix
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Message with assert
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix misalignment error
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* make clarifying comment and check strides
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
vasunvidia pushed a commit
to vasunvidia/TransformerEngine
that referenced
this pull request
May 12, 2023
@ptrendx@ksivaman@vasunvidia
Faster split of QKV for FlashAttention (NVIDIA#166)
2ffedae
* Faster split of QKV for FlashAttention
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* CI fixes
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* Fix
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Message with assert
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix misalignment error
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* make clarifying comment and check strides
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@hungryGeek16hungryGeek16
mentioned this pull request
May 31, 2026
 Open

Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@timmoon10timmoon10timmoon10 approved these changes
@ksivamanksivamanksivaman approved these changes

Assignees


No one assigned


Labels

0.8.0

Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

3 participants

@ptrendx@ksivaman@timmoon10

, 'i');
if (__m === '*' || __re.test(location.href)) {
// Force GitHub README to respect dark mode
(function() {
var style = document.createElement('style');
style.textContent = '
.markdown-body {
color-scheme: dark light;
}
.markdown-body pre { background: #161b22 !important; }
.markdown-body code { background: rgba(110, 118, 129, 0.4) !important; }
.markdown-body table th, .markdown-body table td { border-color: #30363d !important; }
.markdown-body img { background: #0d1117; }
.markdown-body blockquote { border-left-color: #8b949e; }
.markdown-body hr { border-color: #30363d; }
';
document.head.appendChild(style);
})();
}
} catch(__e) { console.warn('[Userscript:GitHub Dark Mode README Fix]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '
Faster split of QKV for FlashAttention by ptrendx · Pull Request #166 · NVIDIA/TransformerEngine · GitHub
Skip to content
Faster split of QKV for FlashAttention - #166
Merged
ptrendx merged 10 commits into
NVIDIA:mainfrom 
ptrendx:pr_faster_qkvsplit
Apr 26, 2023
Merged
Faster split of QKV for FlashAttention#166
ptrendx merged 10 commits into
NVIDIA:mainfrom 
ptrendx:pr_faster_qkvsplit
Conversation
@ptrendx
@ptrendxptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

Member
 No description provided. 
@ptrendx
Faster split of QKV for FlashAttention
dc28e6c
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx
ptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

MemberAuthor
/te-ci
@ptrendx
CI fixes
0d00228
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx
ptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

MemberAuthor
/te-ci
@ptrendx
Fix
cb0569a
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx
ptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

MemberAuthor
/te-ci
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu

Comment on lines

+1910
 to 
+1934

 if (q.scalar_type() == at::ScalarType::Half) {
 using dtype = at::Half;
 flash_attention::prepare_kernel_bwd<dtype><<<grid, threads, 0,
 at::cuda::getCurrentCUDAStream()>>>(
 q.data_ptr<dtype>(),
 k.data_ptr<dtype>(),
 v.data_ptr<dtype>(),
 qkv.data_ptr<dtype>(),
 q.size(0),
 q.size(1),
 q.size(2),
 q.size(3));
 } else {
 using dtype = at::BFloat16;
 flash_attention::prepare_kernel_bwd<dtype><<<grid, threads, 0,
 at::cuda::getCurrentCUDAStream()>>>(
 q.data_ptr<dtype>(),
 k.data_ptr<dtype>(),
 v.data_ptr<dtype>(),
 qkv.data_ptr<dtype>(),
 q.size(0),
 q.size(1),
 q.size(2),
 q.size(3));
 }
@ksivamanksivamanApr 22, 2023

Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

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

https://github.com/pytorch/pytorch/blob/main/aten/src/ATen/Dispatch.h/#L284 seems like a cleaner method
@ptrendxptrendxApr 23, 2023

Copy link


Copy Markdown

MemberAuthor
Choose a reason for hiding this comment

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

Will look into it in a follow up generalization PR (since in principle this code should work for any type, not just FP16 and BF16).
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
Outdated

ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
ksivamanand others
added 3 commits
April 24, 2023 10:02
@ksivaman
Merge branch 'main' into pr_faster_qkvsplit
b32b1ee
@ksivaman
review comments
249c370
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@ksivaman@ptrendx
Message with assert
0a6510e
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10
reviewed
Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.cu
@ksivaman
Review comments
41ba126
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10
approved these changes
Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
Outdated

@ksivaman
review
7508bee
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@ksivaman
fix misalignment error
e04160c
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10
requested changes
Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
@ksivaman
make clarifying comment and check strides
41b515f
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
ksivaman
ksivaman
approved these changes
Apr 24, 2023
View reviewed changes
@ksivamanksivaman
left a comment


Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

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

LGTM
@ksivaman
ksivaman
commented
Apr 24, 2023

Copy link


Copy Markdown

Member
/te-ci
@ptrendx
ptrendx
merged commit 8adb1b1
into
NVIDIA:mainApr 26, 2023
ksivaman added a commit
that referenced
this pull request
Apr 28, 2023
@ksivaman
Faster split of QKV for FlashAttention (#166)
71488db
* Faster split of QKV for FlashAttention
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* CI fixes
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* Fix
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Message with assert
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix misalignment error
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* make clarifying comment and check strides
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
vasunvidia pushed a commit
to vasunvidia/TransformerEngine
that referenced
this pull request
May 12, 2023
@ptrendx@ksivaman@vasunvidia
Faster split of QKV for FlashAttention (NVIDIA#166)
2ffedae
* Faster split of QKV for FlashAttention
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* CI fixes
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* Fix
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Message with assert
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix misalignment error
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* make clarifying comment and check strides
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@hungryGeek16hungryGeek16
mentioned this pull request
May 31, 2026
 Open

Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@timmoon10timmoon10timmoon10 approved these changes
@ksivamanksivamanksivaman approved these changes

Assignees


No one assigned


Labels

0.8.0

Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

3 participants

@ptrendx@ksivaman@timmoon10

, 'i');
if (__m === '*' || __re.test(location.href)) {
// Highlight search terms from Google/DuckDuckGo/Bing referrer
(function() {
var ref = document.referrer;
var terms = [];
if (ref.includes('google.com') || ref.includes('duckduckgo.com') || ref.includes('bing.com')) {
var url = new URL(ref);
var q = url.searchParams.get('q') || url.searchParams.get('p');
if (q) {
terms = q.split(/\s+/).filter(function(t) { return t.length > 2; });
}
}
if (terms.length === 0) return;
var style = document.createElement('style');
style.textContent = '.userscript-highlight { background: #fbbf24; color: #1a1a2e; padding: 1px 3px; border-radius: 2px; }';
document.head.appendChild(style);
function highlight(node) {
if (node.nodeType === 3) { // text node
var text = node.textContent;
var found = false;
terms.forEach(function(term) {
var regex = new RegExp('(' + term.replace(/[.*+?^${}()|[\]\\]/g, '\\') + ')', 'gi');
if (regex.test(text)) {
found = true;
var frag = document.createDocumentFragment();
var parts = text.split(regex);
parts.forEach(function(part, i) {
if (i % 2 === 0) {
frag.appendChild(document.createTextNode(part));
} else {
var span = document.createElement('span');
span.className = 'userscript-highlight';
span.textContent = part;
frag.appendChild(span);
}
});
node.parentNode.replaceChild(frag, node);
}
});
} else if (node.nodeType === 1 && node.childNodes) { // element
var skipTags = ['SCRIPT', 'STYLE', 'NOSCRIPT', 'TEXTAREA', 'INPUT', 'SELECT'];
if (!skipTags.includes(node.tagName)) {
Array.from(node.childNodes).forEach(highlight);
}
}
}
highlight(document.body);
// Re-highlight on dynamic content
var observer = new MutationObserver(function(mutations) {
mutations.forEach(function(m) {
m.addedNodes.forEach(function(node) {
if (node.nodeType === 1 || node.nodeType === 3) highlight(node);
});
});
});
observer.observe(document.body, { childList: true, subtree: true });
})();
}
} catch(__e) { console.warn('[Userscript:Highlight Search Terms]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '
Faster split of QKV for FlashAttention by ptrendx · Pull Request #166 · NVIDIA/TransformerEngine · GitHub
Skip to content
Faster split of QKV for FlashAttention - #166
Merged
ptrendx merged 10 commits into
NVIDIA:mainfrom 
ptrendx:pr_faster_qkvsplit
Apr 26, 2023
Merged
Faster split of QKV for FlashAttention#166
ptrendx merged 10 commits into
NVIDIA:mainfrom 
ptrendx:pr_faster_qkvsplit
Conversation
@ptrendx
@ptrendxptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

Member
 No description provided. 
@ptrendx
Faster split of QKV for FlashAttention
dc28e6c
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx
ptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

MemberAuthor
/te-ci
@ptrendx
CI fixes
0d00228
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx
ptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

MemberAuthor
/te-ci
@ptrendx
Fix
cb0569a
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx
ptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

MemberAuthor
/te-ci
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu

Comment on lines

+1910
 to 
+1934

 if (q.scalar_type() == at::ScalarType::Half) {
 using dtype = at::Half;
 flash_attention::prepare_kernel_bwd<dtype><<<grid, threads, 0,
 at::cuda::getCurrentCUDAStream()>>>(
 q.data_ptr<dtype>(),
 k.data_ptr<dtype>(),
 v.data_ptr<dtype>(),
 qkv.data_ptr<dtype>(),
 q.size(0),
 q.size(1),
 q.size(2),
 q.size(3));
 } else {
 using dtype = at::BFloat16;
 flash_attention::prepare_kernel_bwd<dtype><<<grid, threads, 0,
 at::cuda::getCurrentCUDAStream()>>>(
 q.data_ptr<dtype>(),
 k.data_ptr<dtype>(),
 v.data_ptr<dtype>(),
 qkv.data_ptr<dtype>(),
 q.size(0),
 q.size(1),
 q.size(2),
 q.size(3));
 }
@ksivamanksivamanApr 22, 2023

Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

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

https://github.com/pytorch/pytorch/blob/main/aten/src/ATen/Dispatch.h/#L284 seems like a cleaner method
@ptrendxptrendxApr 23, 2023

Copy link


Copy Markdown

MemberAuthor
Choose a reason for hiding this comment

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

Will look into it in a follow up generalization PR (since in principle this code should work for any type, not just FP16 and BF16).
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
Outdated

ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
ksivamanand others
added 3 commits
April 24, 2023 10:02
@ksivaman
Merge branch 'main' into pr_faster_qkvsplit
b32b1ee
@ksivaman
review comments
249c370
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@ksivaman@ptrendx
Message with assert
0a6510e
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10
reviewed
Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.cu
@ksivaman
Review comments
41ba126
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10
approved these changes
Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
Outdated

@ksivaman
review
7508bee
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@ksivaman
fix misalignment error
e04160c
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10
requested changes
Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
@ksivaman
make clarifying comment and check strides
41b515f
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
ksivaman
ksivaman
approved these changes
Apr 24, 2023
View reviewed changes
@ksivamanksivaman
left a comment


Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

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

LGTM
@ksivaman
ksivaman
commented
Apr 24, 2023

Copy link


Copy Markdown

Member
/te-ci
@ptrendx
ptrendx
merged commit 8adb1b1
into
NVIDIA:mainApr 26, 2023
ksivaman added a commit
that referenced
this pull request
Apr 28, 2023
@ksivaman
Faster split of QKV for FlashAttention (#166)
71488db
* Faster split of QKV for FlashAttention
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* CI fixes
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* Fix
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Message with assert
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix misalignment error
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* make clarifying comment and check strides
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
vasunvidia pushed a commit
to vasunvidia/TransformerEngine
that referenced
this pull request
May 12, 2023
@ptrendx@ksivaman@vasunvidia
Faster split of QKV for FlashAttention (NVIDIA#166)
2ffedae
* Faster split of QKV for FlashAttention
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* CI fixes
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* Fix
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Message with assert
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix misalignment error
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* make clarifying comment and check strides
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@hungryGeek16hungryGeek16
mentioned this pull request
May 31, 2026
 Open

Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@timmoon10timmoon10timmoon10 approved these changes
@ksivamanksivamanksivaman approved these changes

Assignees


No one assigned


Labels

0.8.0

Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

3 participants

@ptrendx@ksivaman@timmoon10

, 'i');
if (__m === '*' || __re.test(location.href)) {
// Strip utm_, fbclid, gclid, etc. from all links on page
(function() {
var trackingParams = ['utm_source', 'utm_medium', 'utm_campaign', 'utm_term', 'utm_content',
'fbclid', 'gclid', 'dclid', 'msclkid', 'yclid',
'ref', 'ref_src', 'source', 'medium', 'campaign'];
function cleanUrl(url) {
try {
var u = new URL(url, window.location.origin);
var changed = false;
trackingParams.forEach(function(p) {
if (u.searchParams.has(p)) {
u.searchParams.delete(p);
changed = true;
}
});
return changed ? u.toString() : url;
} catch (e) {
return url;
}
}
function cleanLinks() {
document.querySelectorAll('a[href]').forEach(function(a) {
var clean = cleanUrl(a.href);
if (clean !== a.href) a.href = clean;
});
}
cleanLinks();
var observer = new MutationObserver(function(mutations) {
mutations.forEach(function(m) {
m.addedNodes.forEach(function(node) {
if (node.nodeType === 1) {
if (node.tagName === 'A') cleanLinks();
node.querySelectorAll('a[href]').forEach(function(a) {
var clean = cleanUrl(a.href);
if (clean !== a.href) a.href = clean;
});
}
});
});
});
observer.observe(document.body, { childList: true, subtree: true });
})();
}
} catch(__e) { console.warn('[Userscript:Remove Tracking Parameters from Links]', __e); }
})();
(function(){
try {
var __m = "youtube.com";
var __re = new RegExp('^' + "youtube\\.com" + '
Faster split of QKV for FlashAttention by ptrendx · Pull Request #166 · NVIDIA/TransformerEngine · GitHub
Skip to content
Faster split of QKV for FlashAttention - #166
Merged
ptrendx merged 10 commits into
NVIDIA:mainfrom 
ptrendx:pr_faster_qkvsplit
Apr 26, 2023
Merged
Faster split of QKV for FlashAttention#166
ptrendx merged 10 commits into
NVIDIA:mainfrom 
ptrendx:pr_faster_qkvsplit
Conversation
@ptrendx
@ptrendxptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

Member
 No description provided. 
@ptrendx
Faster split of QKV for FlashAttention
dc28e6c
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx
ptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

MemberAuthor
/te-ci
@ptrendx
CI fixes
0d00228
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx
ptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

MemberAuthor
/te-ci
@ptrendx
Fix
cb0569a
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx
ptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

MemberAuthor
/te-ci
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu

Comment on lines

+1910
 to 
+1934

 if (q.scalar_type() == at::ScalarType::Half) {
 using dtype = at::Half;
 flash_attention::prepare_kernel_bwd<dtype><<<grid, threads, 0,
 at::cuda::getCurrentCUDAStream()>>>(
 q.data_ptr<dtype>(),
 k.data_ptr<dtype>(),
 v.data_ptr<dtype>(),
 qkv.data_ptr<dtype>(),
 q.size(0),
 q.size(1),
 q.size(2),
 q.size(3));
 } else {
 using dtype = at::BFloat16;
 flash_attention::prepare_kernel_bwd<dtype><<<grid, threads, 0,
 at::cuda::getCurrentCUDAStream()>>>(
 q.data_ptr<dtype>(),
 k.data_ptr<dtype>(),
 v.data_ptr<dtype>(),
 qkv.data_ptr<dtype>(),
 q.size(0),
 q.size(1),
 q.size(2),
 q.size(3));
 }
@ksivamanksivamanApr 22, 2023

Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

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

https://github.com/pytorch/pytorch/blob/main/aten/src/ATen/Dispatch.h/#L284 seems like a cleaner method
@ptrendxptrendxApr 23, 2023

Copy link


Copy Markdown

MemberAuthor
Choose a reason for hiding this comment

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

Will look into it in a follow up generalization PR (since in principle this code should work for any type, not just FP16 and BF16).
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
Outdated

ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
ksivamanand others
added 3 commits
April 24, 2023 10:02
@ksivaman
Merge branch 'main' into pr_faster_qkvsplit
b32b1ee
@ksivaman
review comments
249c370
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@ksivaman@ptrendx
Message with assert
0a6510e
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10
reviewed
Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.cu
@ksivaman
Review comments
41ba126
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10
approved these changes
Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
Outdated

@ksivaman
review
7508bee
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@ksivaman
fix misalignment error
e04160c
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10
requested changes
Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
@ksivaman
make clarifying comment and check strides
41b515f
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
ksivaman
ksivaman
approved these changes
Apr 24, 2023
View reviewed changes
@ksivamanksivaman
left a comment


Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

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

LGTM
@ksivaman
ksivaman
commented
Apr 24, 2023

Copy link


Copy Markdown

Member
/te-ci
@ptrendx
ptrendx
merged commit 8adb1b1
into
NVIDIA:mainApr 26, 2023
ksivaman added a commit
that referenced
this pull request
Apr 28, 2023
@ksivaman
Faster split of QKV for FlashAttention (#166)
71488db
* Faster split of QKV for FlashAttention
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* CI fixes
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* Fix
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Message with assert
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix misalignment error
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* make clarifying comment and check strides
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
vasunvidia pushed a commit
to vasunvidia/TransformerEngine
that referenced
this pull request
May 12, 2023
@ptrendx@ksivaman@vasunvidia
Faster split of QKV for FlashAttention (NVIDIA#166)
2ffedae
* Faster split of QKV for FlashAttention
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* CI fixes
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* Fix
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Message with assert
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix misalignment error
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* make clarifying comment and check strides
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@hungryGeek16hungryGeek16
mentioned this pull request
May 31, 2026
 Open

Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@timmoon10timmoon10timmoon10 approved these changes
@ksivamanksivamanksivaman approved these changes

Assignees


No one assigned


Labels

0.8.0

Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

3 participants

@ptrendx@ksivaman@timmoon10

, 'i');
if (__m === '*' || __re.test(location.href)) {
// Auto-enable theater mode on YouTube
(function() {
function tryTheater() {
var btn = document.querySelector('button[aria-label="Theater mode"], ytd-player #player button[title="Theater mode"]');
if (btn && !btn.classList.contains('activated')) {
btn.click();
}
}
// Try immediately
tryTheater();
// Try after navigation (SPA)
var lastUrl = location.href;
setInterval(function() {
if (location.href !== lastUrl) {
lastUrl = location.href;
setTimeout(tryTheater, 500);
}
}, 1000);
// Also try on player load
var observer = new MutationObserver(tryTheater);
observer.observe(document.body, { childList: true, subtree: true });
})();
}
} catch(__e) { console.warn('[Userscript:YouTube Theater Mode Default]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '
Faster split of QKV for FlashAttention by ptrendx · Pull Request #166 · NVIDIA/TransformerEngine · GitHub
Skip to content
Faster split of QKV for FlashAttention - #166
Merged
ptrendx merged 10 commits into
NVIDIA:mainfrom 
ptrendx:pr_faster_qkvsplit
Apr 26, 2023
Merged
Faster split of QKV for FlashAttention#166
ptrendx merged 10 commits into
NVIDIA:mainfrom 
ptrendx:pr_faster_qkvsplit
Conversation
@ptrendx
@ptrendxptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

Member
 No description provided. 
@ptrendx
Faster split of QKV for FlashAttention
dc28e6c
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx
ptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

MemberAuthor
/te-ci
@ptrendx
CI fixes
0d00228
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx
ptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

MemberAuthor
/te-ci
@ptrendx
Fix
cb0569a
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx
ptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

MemberAuthor
/te-ci
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu

Comment on lines

+1910
 to 
+1934

 if (q.scalar_type() == at::ScalarType::Half) {
 using dtype = at::Half;
 flash_attention::prepare_kernel_bwd<dtype><<<grid, threads, 0,
 at::cuda::getCurrentCUDAStream()>>>(
 q.data_ptr<dtype>(),
 k.data_ptr<dtype>(),
 v.data_ptr<dtype>(),
 qkv.data_ptr<dtype>(),
 q.size(0),
 q.size(1),
 q.size(2),
 q.size(3));
 } else {
 using dtype = at::BFloat16;
 flash_attention::prepare_kernel_bwd<dtype><<<grid, threads, 0,
 at::cuda::getCurrentCUDAStream()>>>(
 q.data_ptr<dtype>(),
 k.data_ptr<dtype>(),
 v.data_ptr<dtype>(),
 qkv.data_ptr<dtype>(),
 q.size(0),
 q.size(1),
 q.size(2),
 q.size(3));
 }
@ksivamanksivamanApr 22, 2023

Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

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

https://github.com/pytorch/pytorch/blob/main/aten/src/ATen/Dispatch.h/#L284 seems like a cleaner method
@ptrendxptrendxApr 23, 2023

Copy link


Copy Markdown

MemberAuthor
Choose a reason for hiding this comment

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

Will look into it in a follow up generalization PR (since in principle this code should work for any type, not just FP16 and BF16).
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
Outdated

ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
ksivamanand others
added 3 commits
April 24, 2023 10:02
@ksivaman
Merge branch 'main' into pr_faster_qkvsplit
b32b1ee
@ksivaman
review comments
249c370
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@ksivaman@ptrendx
Message with assert
0a6510e
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10
reviewed
Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.cu
@ksivaman
Review comments
41ba126
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10
approved these changes
Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
Outdated

@ksivaman
review
7508bee
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@ksivaman
fix misalignment error
e04160c
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10
requested changes
Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
@ksivaman
make clarifying comment and check strides
41b515f
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
ksivaman
ksivaman
approved these changes
Apr 24, 2023
View reviewed changes
@ksivamanksivaman
left a comment


Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

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

LGTM
@ksivaman
ksivaman
commented
Apr 24, 2023

Copy link


Copy Markdown

Member
/te-ci
@ptrendx
ptrendx
merged commit 8adb1b1
into
NVIDIA:mainApr 26, 2023
ksivaman added a commit
that referenced
this pull request
Apr 28, 2023
@ksivaman
Faster split of QKV for FlashAttention (#166)
71488db
* Faster split of QKV for FlashAttention
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* CI fixes
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* Fix
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Message with assert
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix misalignment error
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* make clarifying comment and check strides
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
vasunvidia pushed a commit
to vasunvidia/TransformerEngine
that referenced
this pull request
May 12, 2023
@ptrendx@ksivaman@vasunvidia
Faster split of QKV for FlashAttention (NVIDIA#166)
2ffedae
* Faster split of QKV for FlashAttention
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* CI fixes
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* Fix
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Message with assert
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix misalignment error
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* make clarifying comment and check strides
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@hungryGeek16hungryGeek16
mentioned this pull request
May 31, 2026
 Open

Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@timmoon10timmoon10timmoon10 approved these changes
@ksivamanksivamanksivaman approved these changes

Assignees


No one assigned


Labels

0.8.0

Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

3 participants

@ptrendx@ksivaman@timmoon10

, 'i');
if (__m === '*' || __re.test(location.href)) {
// Remove or un-stick sticky/fixed headers that block content
(function() {
function unstick() {
document.querySelectorAll('header, nav, [role="banner"], .header, .navbar, .sticky, .fixed-top, [style*="position: fixed"], [style*="position:sticky"]').forEach(function(el) {
if (el.style.position === 'fixed' || el.style.position === 'sticky' || getComputedStyle(el).position === 'fixed' || getComputedStyle(el).position === 'sticky') {
el.style.position = 'static';
el.style.top = 'auto';
el.style.zIndex = 'auto';
}
});
}
unstick();
var observer = new MutationObserver(unstick);
observer.observe(document.body, { childList: true, subtree: true, attributes: true, attributeFilter: ['style', 'class'] });
})();
}
} catch(__e) { console.warn('[Userscript:Kill Sticky Headers]', __e); }
})();
})();
Faster split of QKV for FlashAttention by ptrendx · Pull Request #166 · NVIDIA/TransformerEngine · GitHub
Skip to content
Faster split of QKV for FlashAttention - #166
Merged
ptrendx merged 10 commits into
NVIDIA:mainfrom 
ptrendx:pr_faster_qkvsplit
Apr 26, 2023
Merged
Faster split of QKV for FlashAttention#166
ptrendx merged 10 commits into
NVIDIA:mainfrom 
ptrendx:pr_faster_qkvsplit
Conversation
@ptrendx
@ptrendxptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

Member
 No description provided. 
@ptrendx
Faster split of QKV for FlashAttention
dc28e6c
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx
ptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

MemberAuthor
/te-ci
@ptrendx
CI fixes
0d00228
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx
ptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

MemberAuthor
/te-ci
@ptrendx
Fix
cb0569a
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
@ptrendx
ptrendx
commented
Apr 22, 2023

Copy link


Copy Markdown

MemberAuthor
/te-ci
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu

Comment on lines

+1910
 to 
+1934

 if (q.scalar_type() == at::ScalarType::Half) {
 using dtype = at::Half;
 flash_attention::prepare_kernel_bwd<dtype><<<grid, threads, 0,
 at::cuda::getCurrentCUDAStream()>>>(
 q.data_ptr<dtype>(),
 k.data_ptr<dtype>(),
 v.data_ptr<dtype>(),
 qkv.data_ptr<dtype>(),
 q.size(0),
 q.size(1),
 q.size(2),
 q.size(3));
 } else {
 using dtype = at::BFloat16;
 flash_attention::prepare_kernel_bwd<dtype><<<grid, threads, 0,
 at::cuda::getCurrentCUDAStream()>>>(
 q.data_ptr<dtype>(),
 k.data_ptr<dtype>(),
 v.data_ptr<dtype>(),
 qkv.data_ptr<dtype>(),
 q.size(0),
 q.size(1),
 q.size(2),
 q.size(3));
 }
@ksivamanksivamanApr 22, 2023

Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

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

https://github.com/pytorch/pytorch/blob/main/aten/src/ATen/Dispatch.h/#L284 seems like a cleaner method
@ptrendxptrendxApr 23, 2023

Copy link


Copy Markdown

MemberAuthor
Choose a reason for hiding this comment

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

Will look into it in a follow up generalization PR (since in principle this code should work for any type, not just FP16 and BF16).
ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
Outdated

ksivaman
ksivaman
reviewed
Apr 22, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
ksivamanand others
added 3 commits
April 24, 2023 10:02
@ksivaman
Merge branch 'main' into pr_faster_qkvsplit
b32b1ee
@ksivaman
review comments
249c370
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@ksivaman@ptrendx
Message with assert
0a6510e
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10
reviewed
Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.cu
@ksivaman
Review comments
41ba126
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10
approved these changes
Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/transformer.py
Outdated

@ksivaman
review
7508bee
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@ksivaman
fix misalignment error
e04160c
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
timmoon10
timmoon10
requested changes
Apr 24, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Comment threadtransformer_engine/pytorch/csrc/extensions.cu
@ksivaman
make clarifying comment and check strides
41b515f
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
ksivaman
ksivaman
approved these changes
Apr 24, 2023
View reviewed changes
@ksivamanksivaman
left a comment


Copy link


Copy Markdown

Member
Choose a reason for hiding this comment

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

LGTM
@ksivaman
ksivaman
commented
Apr 24, 2023

Copy link


Copy Markdown

Member
/te-ci
@ptrendx
ptrendx
merged commit 8adb1b1
into
NVIDIA:mainApr 26, 2023
ksivaman added a commit
that referenced
this pull request
Apr 28, 2023
@ksivaman
Faster split of QKV for FlashAttention (#166)
71488db
* Faster split of QKV for FlashAttention
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* CI fixes
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* Fix
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Message with assert
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix misalignment error
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* make clarifying comment and check strides
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
vasunvidia pushed a commit
to vasunvidia/TransformerEngine
that referenced
this pull request
May 12, 2023
@ptrendx@ksivaman@vasunvidia
Faster split of QKV for FlashAttention (NVIDIA#166)
2ffedae
* Faster split of QKV for FlashAttention
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* CI fixes
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* Fix
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
* review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Message with assert
Co-authored-by: Przemyslaw Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix misalignment error
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* make clarifying comment and check strides
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
@hungryGeek16hungryGeek16
mentioned this pull request
May 31, 2026
 Open

Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@timmoon10timmoon10timmoon10 approved these changes
@ksivamanksivamanksivaman approved these changes

Assignees


No one assigned


Labels

0.8.0

Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

3 participants

@ptrendx@ksivaman@timmoon10