Skip to content

Increase number of FP8 tensors per GEMM - #22

Merged
ptrendx merged 12 commits into
NVIDIA:mainfrom
vasunvidia:main
Feb 3, 2023
Merged

Increase number of FP8 tensors per GEMM#22
ptrendx merged 12 commits into
NVIDIA:mainfrom
vasunvidia:main

Conversation

@vasunvidia

@vasunvidiavasunvidia commented Oct 28, 2022

Copy link
Copy Markdown
Collaborator

No description provided.

@ptrendx

ptrendx commented Oct 28, 2022

Copy link
Copy Markdown
Member

@vasunvidia Please sign your commits (see CONTRIBUTING.rst)

@ptrendx

ptrendx commented Oct 28, 2022

Copy link
Copy Markdown
Member

/blossom-ci

@vasunvidia

vasunvidia commented Oct 29, 2022

Copy link
Copy Markdown
CollaboratorAuthor

@vasunvidia Please sign your commits (see CONTRIBUTING.rst)

@vasunvidia Please sign your commits (see CONTRIBUTING.rst)

Signed the commit. Thanks.

@ptrendx

ptrendx commented Nov 28, 2022

Copy link
Copy Markdown
Member

@ksivaman Could you review this? Thanks :-)!

vasunvidia
vasunvidia commented Dec 14, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu Outdated
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu Outdated
CUBLASLT_MATMUL_DESC_AMAX_D_POINTER,
&D_amax,
sizeof(D_amax)));
NVTE_CHECK_CUBLAS(cublasLtMatrixLayoutCreate(&Cdesc, bias_type, m, n, ldd));

@vasunvidiavasunvidiaDec 14, 2022

Copy link
Copy Markdown
CollaboratorAuthor

Choose a reason for hiding this comment

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

What if C desc is same as D desc?

@vasunvidiavasunvidiaJan 8, 2023

Copy link
Copy Markdown
CollaboratorAuthor

Choose a reason for hiding this comment

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

For FP8 output, C cannot be FP8. So C type should be same as bias_type for FP8 output type and D_type for other output types.

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu Outdated
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu Outdated
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu Outdated
Comment threadtransformer_engine/pytorch/cpp_extensions.py
Comment threadtransformer_engine/pytorch/cpp_extensions.py Outdated
return_output = True

out_dtype = tex.DType.kFloat32 if fp32_output else TE_DType[out_dtype]
bias_dtype = output_dtype if bias is None else TE_DType[bias.dtype]

@vasunvidiavasunvidiaDec 14, 2022

Copy link
Copy Markdown
CollaboratorAuthor

Choose a reason for hiding this comment

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

SHould not require bias_type in C api.

@vasunvidiavasunvidiaDec 14, 2022

Copy link
Copy Markdown
CollaboratorAuthor

Choose a reason for hiding this comment

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

Check with cublas team

@vasunvidiavasunvidiaJan 8, 2023

Copy link
Copy Markdown
CollaboratorAuthor

Choose a reason for hiding this comment

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

@ksivaman Could you remind why bias_type should not be required?

Comment threadtransformer_engine/pytorch/cpp_extensions.py Outdated
Comment threadtransformer_engine/pytorch/csrc/extensions.h Outdated
Comment threadtransformer_engine/pytorch/csrc/extensions.cu Outdated
ptrendx
ptrendx reviewed Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu Outdated

@ptrendxptrendxDec 16, 2022

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.

You are leaking memory here since Cdesc was already created.

@vasunvidiavasunvidiaJan 8, 2023

Copy link
Copy Markdown
CollaboratorAuthor

Choose a reason for hiding this comment

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

Thanks for pointing it out. Will fix this.

ptrendx
ptrendx reviewed Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu Outdated

@ptrendxptrendxDec 16, 2022

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.

TBH I fail to see why do you have to set C descriptor - it is there for the beta=1 case, right? So it should always be of the same type as D?

@vasunvidiavasunvidiaJan 8, 2023

Copy link
Copy Markdown
CollaboratorAuthor

Choose a reason for hiding this comment

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

For FP8 GEMM output, cublas doesn't support FP8 C_type. So we need to set CDesc to bias type for FP8 output and D type for others. Does it make sense?
image

ptrendx
ptrendx reviewed Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu Outdated

@ptrendxptrendxDec 16, 2022

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.

This breaks beta=1 case that we use for wgrad accumulation, no?

@vasunvidiavasunvidiaJan 8, 2023

Copy link
Copy Markdown
CollaboratorAuthor

Choose a reason for hiding this comment

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

Got it. Will address this.

vasunvidia
vasunvidia commented Jan 17, 2023
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu Outdated
ptrendx
ptrendx reviewed Jan 31, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/ts_fp8_op.cpp Outdated

@ptrendxptrendxJan 31, 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.

Why do you need to clone them?

@vasunvidiavasunvidiaJan 31, 2023

Copy link
Copy Markdown
CollaboratorAuthor

Choose a reason for hiding this comment

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

I just copied the usage for other scale factors such as
at::Tensor A_scale_inverse_arg = A_scale_inverse.clone();

Is it unnecessary?

@ptrendxptrendxJan 31, 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.

@ksivaman could you comment on that?

@ksivamanksivamanJan 31, 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.

Yes this shouldn't be cloned for 2 reasons:

  1. Unnecessary malloc,
  2. and more importantly, we want to populate the original amax tensor received in the arguments, which would remain unchanged here.

@vasunvidiavasunvidiaJan 31, 2023

Copy link
Copy Markdown
CollaboratorAuthor

Choose a reason for hiding this comment

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

Thanks. I'll make the change.

@ptrendxptrendxFeb 1, 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.

@ksivaman That begs the question why are there the other clones for scale inverses? I believe those are unnecessary as well and create additional (albeit pretty small)D2D copies before every FP8 gemm.

@ksivamanksivamanFeb 1, 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.

Yes, I believe all 7 clones in transformer_engine/pytorch/csrc/ts_fp8_op.cpp can be removed. The scale inverse clones also seem unused? @asfiyab-nvidia Could you comment on why these were added to begin with?

@asfiyab-nvidiaasfiyab-nvidiaFeb 1, 2023

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

@ksivaman These aren't necessary. I'll create a PR with a fix shortly. Thanks for pointing it out

ptrendx
ptrendx reviewed Jan 31, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/module.py Outdated
@ptrendx

ptrendx commented Jan 31, 2023

Copy link
Copy Markdown
Member

/te-ci

@vasunvidia
Increase number of FP8 tensors per GEMM
95910d8
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Enable FP8 output tensor for fp8_gemm
fcbce36
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
[BERT FP8] Initial TE review comments
8a336ab
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Temporary fix for cuda graph non convergence
320182a
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Address review comments-2
ffe4852
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Review comments-3
075e287
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Cleanup
784fb47
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Change for New API
0ccb2e5
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Remove unnecessary clone for D_scale, D_amax
397a1e4
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Avoid Roll for AMAX history size = 1
a575b5e
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Update onnx_te_gemm API
2dad360
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia

vasunvidia commented Feb 2, 2023

Copy link
Copy Markdown
CollaboratorAuthor

/te-ci

1 similar comment
@ptrendx

ptrendx commented Feb 2, 2023

Copy link
Copy Markdown
Member

/te-ci

@vasunvidia
Fix Lint errors
66b98e3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
ptrendx
ptrendx approved these changes Feb 3, 2023
View reviewed changes

@ptrendxptrendx 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, thanks :-).

@ptrendx
ptrendx merged commit 14198f2 into NVIDIA:mainFeb 3, 2023
cyanguwa pushed a commit to cyanguwa/TransformerEngine that referenced this pull request Feb 13, 2023
@vasunvidia@cyanguwa
Increase number of FP8 tensors per GEMM (NVIDIA#22)
f071530
* Increase number of FP8 tensors per GEMM
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Enable FP8 output tensor for fp8_gemm
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* [BERT FP8] Initial TE review comments
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Temporary fix for cuda graph non convergence
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Address review comments-2
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Review comments-3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Cleanup
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Change for New API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Remove unnecessary clone for D_scale, D_amax
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Avoid Roll for AMAX history size = 1
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Update onnx_te_gemm API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Fix Lint errors
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
---------
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
ksivaman added a commit that referenced this pull request Feb 22, 2023
@cyanguwa@ksivaman
flash-attn integration (#62)
79a9fe2
* add flash attention to TransformerLayer
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Add docs for FP8 calibration (#61)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Fix the integer overflow in fused softmax (#60)
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* prefix flash attn env var with NVTE_
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Address steady memory increase and bloated checkpoints (#63)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix env var logic
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix flash attn env var logic again
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* remove d2d copies (#64)
* remove d2d copies
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* cleanup
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Increase number of FP8 tensors per GEMM (#22)
* Increase number of FP8 tensors per GEMM
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Enable FP8 output tensor for fp8_gemm
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* [BERT FP8] Initial TE review comments
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Temporary fix for cuda graph non convergence
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Address review comments-2
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Review comments-3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Cleanup
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Change for New API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Remove unnecessary clone for D_scale, D_amax
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Avoid Roll for AMAX history size = 1
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Update onnx_te_gemm API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Fix Lint errors
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
---------
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Bug fixes from PR 22 (#65)
* Bug fixes from PR 22
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add FP8 tests to ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* bundle unittests for ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* replace rearrange with transpose
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* QKV parameters unfused path fixes and optimization (#66)
* Bug fixes from PR 22
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add FP8 tests to ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Better QKV parameter fusion
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* small fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* keep original param for unfused case to retain externally set attrs
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Fix ONNX exports
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* improve arg naming
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* No need to set data pointers
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Assert memory loc in NoopCat
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Handle case of different memory in param and buffer
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix assert always true
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Reassign params memory to avoid more concats
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Fix gradients when using AMP (#70)
retain grad related attrs while casting
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix pylint violations fixed pyline violations such as trailing white spaces and too long lines Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* fix pylint violation on line 264 with R1719
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* fix two more pylint violations
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* DotProductAttention API
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add docs for attention
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix assert always true
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* check for correct flash-attn version
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* address review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint+build fixes, correct settings for default flash-attn
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* correct version
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review comments and fixes
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix onnx and disable flash-attn export test
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* remove einops dependency
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* cleanup internal API; rm duplication
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* only install TE wheel (exclude flash-attn to rm conflicts)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* forgot to change install wheel path
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* next round review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix flash_attn output
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix QK layer scaling
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* update docs
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review comments and fixes to selective checkpointing
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Charlene Yang <charleney@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Co-authored-by: Charlene Yang <charleney@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
shifangx pushed a commit to shifangx/TransformerEngine that referenced this pull request Feb 10, 2026
@lxd-cumt
Remove use_gems context and call flag_gems.xxx directly (NVIDIA#22)
832a797
- Remove the flag_gems.use_gems() context to avoid context-switching
overhead
- Call flag_gems.xxx directly wherever possible.
Sign up for freeto join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@ptrendxptrendxptrendx approved these changes
@ksivamanksivamanAwaiting requested review from ksivamanksivaman is a code owner
+1 more reviewer
@asfiyab-nvidiaasfiyab-nvidiaasfiyab-nvidia left review comments
Reviewers whose approvals may not affect merge requirements

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

4 participants

@vasunvidia@ptrendx@ksivaman@asfiyab-nvidia
, '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" + '
Increase number of FP8 tensors per GEMM by vasunvidia · Pull Request #22 · NVIDIA/TransformerEngine · GitHub
Skip to content
Increase number of FP8 tensors per GEMM - #22
Merged
ptrendx merged 12 commits into
NVIDIA:mainfrom 
vasunvidia:main
Feb 3, 2023
Merged
Increase number of FP8 tensors per GEMM#22
ptrendx merged 12 commits into
NVIDIA:mainfrom 
vasunvidia:main
Conversation
@vasunvidia
@vasunvidiavasunvidia
commented
Oct 28, 2022

Copy link


Copy Markdown

Collaborator
 No description provided. 
@ptrendx
ptrendx
commented
Oct 28, 2022

Copy link


Copy Markdown

Member
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
@ptrendx
ptrendx
commented
Oct 28, 2022

Copy link


Copy Markdown

Member
/blossom-ci
@vasunvidia
vasunvidia
commented
Oct 29, 2022

Copy link


Copy Markdown

CollaboratorAuthor
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
Signed the commit. Thanks.
@ptrendx
ptrendx
commented
Nov 28, 2022

Copy link


Copy Markdown

Member
@ksivaman Could you review this? Thanks :-)!
vasunvidia
vasunvidia
commented
Dec 14, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

 CUBLASLT_MATMUL_DESC_AMAX_D_POINTER,
 &D_amax,
 sizeof(D_amax)));
 NVTE_CHECK_CUBLAS(cublasLtMatrixLayoutCreate(&Cdesc, bias_type, m, n, ldd));
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

What if C desc is same as D desc?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

For FP8 output, C cannot be FP8. So C type should be same as bias_type for FP8 output type and D_type for other output types.
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/pytorch/cpp_extensions.py
Comment threadtransformer_engine/pytorch/cpp_extensions.py
Outdated

 return_output = True

 out_dtype = tex.DType.kFloat32 if fp32_output else TE_DType[out_dtype]
 bias_dtype = output_dtype if bias is None else TE_DType[bias.dtype]
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

SHould not require bias_type in C api.
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Check with cublas team
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

@ksivaman Could you remind why bias_type should not be required?
Comment threadtransformer_engine/pytorch/cpp_extensions.py
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.h
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

You are leaking memory here since Cdesc was already created.
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Thanks for pointing it out. Will fix this.
ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

TBH I fail to see why do you have to set C descriptor - it is there for the beta=1 case, right? So it should always be of the same type as D?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

For FP8 GEMM output, cublas doesn't support FP8 C_type. So we need to set CDesc to bias type for FP8 output and D type for others. Does it make sense?
image
ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

This breaks beta=1 case that we use for wgrad accumulation, no?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Got it. Will address this.
vasunvidia
vasunvidia
commented
Jan 17, 2023
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

ptrendx
ptrendx
reviewed
Jan 31, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/ts_fp8_op.cpp
Outdated

@ptrendxptrendxJan 31, 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.

Why do you need to clone them?
@vasunvidiavasunvidiaJan 31, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

I just copied the usage for other scale factors such as

at::Tensor A_scale_inverse_arg = A_scale_inverse.clone();
Is it unnecessary?
@ptrendxptrendxJan 31, 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.

@ksivaman could you comment on that?
@ksivamanksivamanJan 31, 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.

Yes this shouldn't be cloned for 2 reasons:
  1. Unnecessary malloc,
  2. and more importantly, we want to populate the original amax tensor received in the arguments, which would remain unchanged here.
@vasunvidiavasunvidiaJan 31, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Thanks. I'll make the change.
@ptrendxptrendxFeb 1, 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.

@ksivaman That begs the question why are there the other clones for scale inverses? I believe those are unnecessary as well and create additional (albeit pretty small)D2D copies before every FP8 gemm.
@ksivamanksivamanFeb 1, 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.

Yes, I believe all 7 clones in transformer_engine/pytorch/csrc/ts_fp8_op.cpp can be removed. The scale inverse clones also seem unused? @asfiyab-nvidia Could you comment on why these were added to begin with?
@asfiyab-nvidiaasfiyab-nvidiaFeb 1, 2023

Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

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

@ksivaman These aren't necessary. I'll create a PR with a fix shortly. Thanks for pointing it out
ptrendx
ptrendx
reviewed
Jan 31, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/module.py
Outdated

@ptrendx
ptrendx
commented
Jan 31, 2023

Copy link


Copy Markdown

Member
/te-ci
@vasunvidia
Increase number of FP8 tensors per GEMM
95910d8
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Enable FP8 output tensor for fp8_gemm
fcbce36
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
[BERT FP8] Initial TE review comments
8a336ab
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Temporary fix for cuda graph non convergence
320182a
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Address review comments-2
ffe4852
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Review comments-3
075e287
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Cleanup
784fb47
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Change for New API
0ccb2e5
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Remove unnecessary clone for D_scale, D_amax
397a1e4
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Avoid Roll for AMAX history size = 1
a575b5e
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Update onnx_te_gemm API
2dad360
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
vasunvidia
commented
Feb 2, 2023

Copy link


Copy Markdown

CollaboratorAuthor
/te-ci

1 similar comment

@ptrendx
ptrendx
commented
Feb 2, 2023

Copy link


Copy Markdown

Member
/te-ci
@vasunvidia
Fix Lint errors
66b98e3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
ptrendx
ptrendx
approved these changes
Feb 3, 2023
View reviewed changes
@ptrendxptrendx
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, thanks :-).
@ptrendx
ptrendx
merged commit 14198f2
into
NVIDIA:mainFeb 3, 2023
cyanguwa pushed a commit
to cyanguwa/TransformerEngine
that referenced
this pull request
Feb 13, 2023
@vasunvidia@cyanguwa
Increase number of FP8 tensors per GEMM (NVIDIA#22)
f071530
* Increase number of FP8 tensors per GEMM
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Enable FP8 output tensor for fp8_gemm
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* [BERT FP8] Initial TE review comments
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Temporary fix for cuda graph non convergence
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Address review comments-2
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Review comments-3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Cleanup
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Change for New API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Remove unnecessary clone for D_scale, D_amax
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Avoid Roll for AMAX history size = 1
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Update onnx_te_gemm API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Fix Lint errors
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
---------
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
ksivaman added a commit
that referenced
this pull request
Feb 22, 2023
@cyanguwa@ksivaman
flash-attn integration (#62)
79a9fe2
* add flash attention to TransformerLayer
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Add docs for FP8 calibration (#61)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Fix the integer overflow in fused softmax (#60)
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* prefix flash attn env var with NVTE_
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Address steady memory increase and bloated checkpoints (#63)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix env var logic
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix flash attn env var logic again
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* remove d2d copies (#64)
* remove d2d copies
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* cleanup
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Increase number of FP8 tensors per GEMM (#22)
* Increase number of FP8 tensors per GEMM
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Enable FP8 output tensor for fp8_gemm
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* [BERT FP8] Initial TE review comments
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Temporary fix for cuda graph non convergence
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Address review comments-2
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Review comments-3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Cleanup
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Change for New API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Remove unnecessary clone for D_scale, D_amax
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Avoid Roll for AMAX history size = 1
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Update onnx_te_gemm API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Fix Lint errors
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
---------
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Bug fixes from PR 22 (#65)
* Bug fixes from PR 22
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add FP8 tests to ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* bundle unittests for ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* replace rearrange with transpose
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* QKV parameters unfused path fixes and optimization (#66)
* Bug fixes from PR 22
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add FP8 tests to ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Better QKV parameter fusion
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* small fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* keep original param for unfused case to retain externally set attrs
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Fix ONNX exports
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* improve arg naming
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* No need to set data pointers
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Assert memory loc in NoopCat
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Handle case of different memory in param and buffer
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix assert always true
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Reassign params memory to avoid more concats
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Fix gradients when using AMP (#70)
retain grad related attrs while casting
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix pylint violations fixed pyline violations such as trailing white spaces and too long lines Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* fix pylint violation on line 264 with R1719
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* fix two more pylint violations
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* DotProductAttention API
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add docs for attention
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix assert always true
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* check for correct flash-attn version
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* address review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint+build fixes, correct settings for default flash-attn
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* correct version
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review comments and fixes
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix onnx and disable flash-attn export test
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* remove einops dependency
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* cleanup internal API; rm duplication
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* only install TE wheel (exclude flash-attn to rm conflicts)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* forgot to change install wheel path
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* next round review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix flash_attn output
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix QK layer scaling
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* update docs
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review comments and fixes to selective checkpointing
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Charlene Yang <charleney@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Co-authored-by: Charlene Yang <charleney@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
shifangx pushed a commit
to shifangx/TransformerEngine
that referenced
this pull request
Feb 10, 2026
@lxd-cumt
Remove use_gems context and call flag_gems.xxx directly (NVIDIA#22)
832a797
- Remove the flag_gems.use_gems() context to avoid context-switching
overhead
- Call flag_gems.xxx directly wherever possible.
Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@ptrendxptrendxptrendx approved these changes
@ksivamanksivamanAwaiting requested review from ksivamanksivaman is a code owner

+1 more reviewer

@asfiyab-nvidiaasfiyab-nvidiaasfiyab-nvidia left review comments
Reviewers whose approvals may not affect merge requirements

Assignees


No one assigned


Labels


None yet


Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

4 participants

@vasunvidia@ptrendx@ksivaman@asfiyab-nvidia

, '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('^' + ".*" + '
Increase number of FP8 tensors per GEMM by vasunvidia · Pull Request #22 · NVIDIA/TransformerEngine · GitHub
Skip to content
Increase number of FP8 tensors per GEMM - #22
Merged
ptrendx merged 12 commits into
NVIDIA:mainfrom 
vasunvidia:main
Feb 3, 2023
Merged
Increase number of FP8 tensors per GEMM#22
ptrendx merged 12 commits into
NVIDIA:mainfrom 
vasunvidia:main
Conversation
@vasunvidia
@vasunvidiavasunvidia
commented
Oct 28, 2022

Copy link


Copy Markdown

Collaborator
 No description provided. 
@ptrendx
ptrendx
commented
Oct 28, 2022

Copy link


Copy Markdown

Member
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
@ptrendx
ptrendx
commented
Oct 28, 2022

Copy link


Copy Markdown

Member
/blossom-ci
@vasunvidia
vasunvidia
commented
Oct 29, 2022

Copy link


Copy Markdown

CollaboratorAuthor
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
Signed the commit. Thanks.
@ptrendx
ptrendx
commented
Nov 28, 2022

Copy link


Copy Markdown

Member
@ksivaman Could you review this? Thanks :-)!
vasunvidia
vasunvidia
commented
Dec 14, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

 CUBLASLT_MATMUL_DESC_AMAX_D_POINTER,
 &D_amax,
 sizeof(D_amax)));
 NVTE_CHECK_CUBLAS(cublasLtMatrixLayoutCreate(&Cdesc, bias_type, m, n, ldd));
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

What if C desc is same as D desc?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

For FP8 output, C cannot be FP8. So C type should be same as bias_type for FP8 output type and D_type for other output types.
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/pytorch/cpp_extensions.py
Comment threadtransformer_engine/pytorch/cpp_extensions.py
Outdated

 return_output = True

 out_dtype = tex.DType.kFloat32 if fp32_output else TE_DType[out_dtype]
 bias_dtype = output_dtype if bias is None else TE_DType[bias.dtype]
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

SHould not require bias_type in C api.
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Check with cublas team
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

@ksivaman Could you remind why bias_type should not be required?
Comment threadtransformer_engine/pytorch/cpp_extensions.py
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.h
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

You are leaking memory here since Cdesc was already created.
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Thanks for pointing it out. Will fix this.
ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

TBH I fail to see why do you have to set C descriptor - it is there for the beta=1 case, right? So it should always be of the same type as D?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

For FP8 GEMM output, cublas doesn't support FP8 C_type. So we need to set CDesc to bias type for FP8 output and D type for others. Does it make sense?
image
ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

This breaks beta=1 case that we use for wgrad accumulation, no?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Got it. Will address this.
vasunvidia
vasunvidia
commented
Jan 17, 2023
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

ptrendx
ptrendx
reviewed
Jan 31, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/ts_fp8_op.cpp
Outdated

@ptrendxptrendxJan 31, 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.

Why do you need to clone them?
@vasunvidiavasunvidiaJan 31, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

I just copied the usage for other scale factors such as

at::Tensor A_scale_inverse_arg = A_scale_inverse.clone();
Is it unnecessary?
@ptrendxptrendxJan 31, 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.

@ksivaman could you comment on that?
@ksivamanksivamanJan 31, 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.

Yes this shouldn't be cloned for 2 reasons:
  1. Unnecessary malloc,
  2. and more importantly, we want to populate the original amax tensor received in the arguments, which would remain unchanged here.
@vasunvidiavasunvidiaJan 31, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Thanks. I'll make the change.
@ptrendxptrendxFeb 1, 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.

@ksivaman That begs the question why are there the other clones for scale inverses? I believe those are unnecessary as well and create additional (albeit pretty small)D2D copies before every FP8 gemm.
@ksivamanksivamanFeb 1, 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.

Yes, I believe all 7 clones in transformer_engine/pytorch/csrc/ts_fp8_op.cpp can be removed. The scale inverse clones also seem unused? @asfiyab-nvidia Could you comment on why these were added to begin with?
@asfiyab-nvidiaasfiyab-nvidiaFeb 1, 2023

Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

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

@ksivaman These aren't necessary. I'll create a PR with a fix shortly. Thanks for pointing it out
ptrendx
ptrendx
reviewed
Jan 31, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/module.py
Outdated

@ptrendx
ptrendx
commented
Jan 31, 2023

Copy link


Copy Markdown

Member
/te-ci
@vasunvidia
Increase number of FP8 tensors per GEMM
95910d8
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Enable FP8 output tensor for fp8_gemm
fcbce36
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
[BERT FP8] Initial TE review comments
8a336ab
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Temporary fix for cuda graph non convergence
320182a
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Address review comments-2
ffe4852
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Review comments-3
075e287
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Cleanup
784fb47
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Change for New API
0ccb2e5
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Remove unnecessary clone for D_scale, D_amax
397a1e4
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Avoid Roll for AMAX history size = 1
a575b5e
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Update onnx_te_gemm API
2dad360
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
vasunvidia
commented
Feb 2, 2023

Copy link


Copy Markdown

CollaboratorAuthor
/te-ci

1 similar comment

@ptrendx
ptrendx
commented
Feb 2, 2023

Copy link


Copy Markdown

Member
/te-ci
@vasunvidia
Fix Lint errors
66b98e3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
ptrendx
ptrendx
approved these changes
Feb 3, 2023
View reviewed changes
@ptrendxptrendx
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, thanks :-).
@ptrendx
ptrendx
merged commit 14198f2
into
NVIDIA:mainFeb 3, 2023
cyanguwa pushed a commit
to cyanguwa/TransformerEngine
that referenced
this pull request
Feb 13, 2023
@vasunvidia@cyanguwa
Increase number of FP8 tensors per GEMM (NVIDIA#22)
f071530
* Increase number of FP8 tensors per GEMM
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Enable FP8 output tensor for fp8_gemm
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* [BERT FP8] Initial TE review comments
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Temporary fix for cuda graph non convergence
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Address review comments-2
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Review comments-3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Cleanup
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Change for New API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Remove unnecessary clone for D_scale, D_amax
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Avoid Roll for AMAX history size = 1
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Update onnx_te_gemm API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Fix Lint errors
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
---------
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
ksivaman added a commit
that referenced
this pull request
Feb 22, 2023
@cyanguwa@ksivaman
flash-attn integration (#62)
79a9fe2
* add flash attention to TransformerLayer
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Add docs for FP8 calibration (#61)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Fix the integer overflow in fused softmax (#60)
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* prefix flash attn env var with NVTE_
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Address steady memory increase and bloated checkpoints (#63)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix env var logic
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix flash attn env var logic again
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* remove d2d copies (#64)
* remove d2d copies
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* cleanup
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Increase number of FP8 tensors per GEMM (#22)
* Increase number of FP8 tensors per GEMM
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Enable FP8 output tensor for fp8_gemm
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* [BERT FP8] Initial TE review comments
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Temporary fix for cuda graph non convergence
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Address review comments-2
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Review comments-3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Cleanup
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Change for New API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Remove unnecessary clone for D_scale, D_amax
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Avoid Roll for AMAX history size = 1
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Update onnx_te_gemm API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Fix Lint errors
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
---------
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Bug fixes from PR 22 (#65)
* Bug fixes from PR 22
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add FP8 tests to ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* bundle unittests for ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* replace rearrange with transpose
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* QKV parameters unfused path fixes and optimization (#66)
* Bug fixes from PR 22
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add FP8 tests to ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Better QKV parameter fusion
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* small fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* keep original param for unfused case to retain externally set attrs
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Fix ONNX exports
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* improve arg naming
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* No need to set data pointers
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Assert memory loc in NoopCat
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Handle case of different memory in param and buffer
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix assert always true
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Reassign params memory to avoid more concats
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Fix gradients when using AMP (#70)
retain grad related attrs while casting
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix pylint violations fixed pyline violations such as trailing white spaces and too long lines Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* fix pylint violation on line 264 with R1719
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* fix two more pylint violations
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* DotProductAttention API
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add docs for attention
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix assert always true
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* check for correct flash-attn version
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* address review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint+build fixes, correct settings for default flash-attn
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* correct version
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review comments and fixes
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix onnx and disable flash-attn export test
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* remove einops dependency
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* cleanup internal API; rm duplication
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* only install TE wheel (exclude flash-attn to rm conflicts)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* forgot to change install wheel path
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* next round review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix flash_attn output
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix QK layer scaling
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* update docs
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review comments and fixes to selective checkpointing
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Charlene Yang <charleney@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Co-authored-by: Charlene Yang <charleney@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
shifangx pushed a commit
to shifangx/TransformerEngine
that referenced
this pull request
Feb 10, 2026
@lxd-cumt
Remove use_gems context and call flag_gems.xxx directly (NVIDIA#22)
832a797
- Remove the flag_gems.use_gems() context to avoid context-switching
overhead
- Call flag_gems.xxx directly wherever possible.
Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@ptrendxptrendxptrendx approved these changes
@ksivamanksivamanAwaiting requested review from ksivamanksivaman is a code owner

+1 more reviewer

@asfiyab-nvidiaasfiyab-nvidiaasfiyab-nvidia left review comments
Reviewers whose approvals may not affect merge requirements

Assignees


No one assigned


Labels


None yet


Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

4 participants

@vasunvidia@ptrendx@ksivaman@asfiyab-nvidia

, '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('^' + ".*" + '
Increase number of FP8 tensors per GEMM by vasunvidia · Pull Request #22 · NVIDIA/TransformerEngine · GitHub
Skip to content
Increase number of FP8 tensors per GEMM - #22
Merged
ptrendx merged 12 commits into
NVIDIA:mainfrom 
vasunvidia:main
Feb 3, 2023
Merged
Increase number of FP8 tensors per GEMM#22
ptrendx merged 12 commits into
NVIDIA:mainfrom 
vasunvidia:main
Conversation
@vasunvidia
@vasunvidiavasunvidia
commented
Oct 28, 2022

Copy link


Copy Markdown

Collaborator
 No description provided. 
@ptrendx
ptrendx
commented
Oct 28, 2022

Copy link


Copy Markdown

Member
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
@ptrendx
ptrendx
commented
Oct 28, 2022

Copy link


Copy Markdown

Member
/blossom-ci
@vasunvidia
vasunvidia
commented
Oct 29, 2022

Copy link


Copy Markdown

CollaboratorAuthor
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
Signed the commit. Thanks.
@ptrendx
ptrendx
commented
Nov 28, 2022

Copy link


Copy Markdown

Member
@ksivaman Could you review this? Thanks :-)!
vasunvidia
vasunvidia
commented
Dec 14, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

 CUBLASLT_MATMUL_DESC_AMAX_D_POINTER,
 &D_amax,
 sizeof(D_amax)));
 NVTE_CHECK_CUBLAS(cublasLtMatrixLayoutCreate(&Cdesc, bias_type, m, n, ldd));
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

What if C desc is same as D desc?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

For FP8 output, C cannot be FP8. So C type should be same as bias_type for FP8 output type and D_type for other output types.
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/pytorch/cpp_extensions.py
Comment threadtransformer_engine/pytorch/cpp_extensions.py
Outdated

 return_output = True

 out_dtype = tex.DType.kFloat32 if fp32_output else TE_DType[out_dtype]
 bias_dtype = output_dtype if bias is None else TE_DType[bias.dtype]
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

SHould not require bias_type in C api.
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Check with cublas team
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

@ksivaman Could you remind why bias_type should not be required?
Comment threadtransformer_engine/pytorch/cpp_extensions.py
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.h
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

You are leaking memory here since Cdesc was already created.
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Thanks for pointing it out. Will fix this.
ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

TBH I fail to see why do you have to set C descriptor - it is there for the beta=1 case, right? So it should always be of the same type as D?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

For FP8 GEMM output, cublas doesn't support FP8 C_type. So we need to set CDesc to bias type for FP8 output and D type for others. Does it make sense?
image
ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

This breaks beta=1 case that we use for wgrad accumulation, no?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Got it. Will address this.
vasunvidia
vasunvidia
commented
Jan 17, 2023
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

ptrendx
ptrendx
reviewed
Jan 31, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/ts_fp8_op.cpp
Outdated

@ptrendxptrendxJan 31, 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.

Why do you need to clone them?
@vasunvidiavasunvidiaJan 31, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

I just copied the usage for other scale factors such as

at::Tensor A_scale_inverse_arg = A_scale_inverse.clone();
Is it unnecessary?
@ptrendxptrendxJan 31, 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.

@ksivaman could you comment on that?
@ksivamanksivamanJan 31, 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.

Yes this shouldn't be cloned for 2 reasons:
  1. Unnecessary malloc,
  2. and more importantly, we want to populate the original amax tensor received in the arguments, which would remain unchanged here.
@vasunvidiavasunvidiaJan 31, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Thanks. I'll make the change.
@ptrendxptrendxFeb 1, 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.

@ksivaman That begs the question why are there the other clones for scale inverses? I believe those are unnecessary as well and create additional (albeit pretty small)D2D copies before every FP8 gemm.
@ksivamanksivamanFeb 1, 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.

Yes, I believe all 7 clones in transformer_engine/pytorch/csrc/ts_fp8_op.cpp can be removed. The scale inverse clones also seem unused? @asfiyab-nvidia Could you comment on why these were added to begin with?
@asfiyab-nvidiaasfiyab-nvidiaFeb 1, 2023

Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

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

@ksivaman These aren't necessary. I'll create a PR with a fix shortly. Thanks for pointing it out
ptrendx
ptrendx
reviewed
Jan 31, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/module.py
Outdated

@ptrendx
ptrendx
commented
Jan 31, 2023

Copy link


Copy Markdown

Member
/te-ci
@vasunvidia
Increase number of FP8 tensors per GEMM
95910d8
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Enable FP8 output tensor for fp8_gemm
fcbce36
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
[BERT FP8] Initial TE review comments
8a336ab
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Temporary fix for cuda graph non convergence
320182a
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Address review comments-2
ffe4852
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Review comments-3
075e287
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Cleanup
784fb47
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Change for New API
0ccb2e5
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Remove unnecessary clone for D_scale, D_amax
397a1e4
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Avoid Roll for AMAX history size = 1
a575b5e
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Update onnx_te_gemm API
2dad360
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
vasunvidia
commented
Feb 2, 2023

Copy link


Copy Markdown

CollaboratorAuthor
/te-ci

1 similar comment

@ptrendx
ptrendx
commented
Feb 2, 2023

Copy link


Copy Markdown

Member
/te-ci
@vasunvidia
Fix Lint errors
66b98e3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
ptrendx
ptrendx
approved these changes
Feb 3, 2023
View reviewed changes
@ptrendxptrendx
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, thanks :-).
@ptrendx
ptrendx
merged commit 14198f2
into
NVIDIA:mainFeb 3, 2023
cyanguwa pushed a commit
to cyanguwa/TransformerEngine
that referenced
this pull request
Feb 13, 2023
@vasunvidia@cyanguwa
Increase number of FP8 tensors per GEMM (NVIDIA#22)
f071530
* Increase number of FP8 tensors per GEMM
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Enable FP8 output tensor for fp8_gemm
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* [BERT FP8] Initial TE review comments
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Temporary fix for cuda graph non convergence
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Address review comments-2
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Review comments-3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Cleanup
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Change for New API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Remove unnecessary clone for D_scale, D_amax
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Avoid Roll for AMAX history size = 1
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Update onnx_te_gemm API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Fix Lint errors
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
---------
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
ksivaman added a commit
that referenced
this pull request
Feb 22, 2023
@cyanguwa@ksivaman
flash-attn integration (#62)
79a9fe2
* add flash attention to TransformerLayer
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Add docs for FP8 calibration (#61)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Fix the integer overflow in fused softmax (#60)
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* prefix flash attn env var with NVTE_
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Address steady memory increase and bloated checkpoints (#63)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix env var logic
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix flash attn env var logic again
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* remove d2d copies (#64)
* remove d2d copies
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* cleanup
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Increase number of FP8 tensors per GEMM (#22)
* Increase number of FP8 tensors per GEMM
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Enable FP8 output tensor for fp8_gemm
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* [BERT FP8] Initial TE review comments
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Temporary fix for cuda graph non convergence
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Address review comments-2
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Review comments-3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Cleanup
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Change for New API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Remove unnecessary clone for D_scale, D_amax
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Avoid Roll for AMAX history size = 1
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Update onnx_te_gemm API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Fix Lint errors
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
---------
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Bug fixes from PR 22 (#65)
* Bug fixes from PR 22
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add FP8 tests to ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* bundle unittests for ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* replace rearrange with transpose
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* QKV parameters unfused path fixes and optimization (#66)
* Bug fixes from PR 22
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add FP8 tests to ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Better QKV parameter fusion
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* small fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* keep original param for unfused case to retain externally set attrs
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Fix ONNX exports
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* improve arg naming
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* No need to set data pointers
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Assert memory loc in NoopCat
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Handle case of different memory in param and buffer
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix assert always true
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Reassign params memory to avoid more concats
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Fix gradients when using AMP (#70)
retain grad related attrs while casting
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix pylint violations fixed pyline violations such as trailing white spaces and too long lines Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* fix pylint violation on line 264 with R1719
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* fix two more pylint violations
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* DotProductAttention API
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add docs for attention
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix assert always true
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* check for correct flash-attn version
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* address review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint+build fixes, correct settings for default flash-attn
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* correct version
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review comments and fixes
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix onnx and disable flash-attn export test
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* remove einops dependency
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* cleanup internal API; rm duplication
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* only install TE wheel (exclude flash-attn to rm conflicts)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* forgot to change install wheel path
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* next round review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix flash_attn output
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix QK layer scaling
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* update docs
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review comments and fixes to selective checkpointing
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Charlene Yang <charleney@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Co-authored-by: Charlene Yang <charleney@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
shifangx pushed a commit
to shifangx/TransformerEngine
that referenced
this pull request
Feb 10, 2026
@lxd-cumt
Remove use_gems context and call flag_gems.xxx directly (NVIDIA#22)
832a797
- Remove the flag_gems.use_gems() context to avoid context-switching
overhead
- Call flag_gems.xxx directly wherever possible.
Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@ptrendxptrendxptrendx approved these changes
@ksivamanksivamanAwaiting requested review from ksivamanksivaman is a code owner

+1 more reviewer

@asfiyab-nvidiaasfiyab-nvidiaasfiyab-nvidia left review comments
Reviewers whose approvals may not affect merge requirements

Assignees


No one assigned


Labels


None yet


Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

4 participants

@vasunvidia@ptrendx@ksivaman@asfiyab-nvidia

, '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" + '
Increase number of FP8 tensors per GEMM by vasunvidia · Pull Request #22 · NVIDIA/TransformerEngine · GitHub
Skip to content
Increase number of FP8 tensors per GEMM - #22
Merged
ptrendx merged 12 commits into
NVIDIA:mainfrom 
vasunvidia:main
Feb 3, 2023
Merged
Increase number of FP8 tensors per GEMM#22
ptrendx merged 12 commits into
NVIDIA:mainfrom 
vasunvidia:main
Conversation
@vasunvidia
@vasunvidiavasunvidia
commented
Oct 28, 2022

Copy link


Copy Markdown

Collaborator
 No description provided. 
@ptrendx
ptrendx
commented
Oct 28, 2022

Copy link


Copy Markdown

Member
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
@ptrendx
ptrendx
commented
Oct 28, 2022

Copy link


Copy Markdown

Member
/blossom-ci
@vasunvidia
vasunvidia
commented
Oct 29, 2022

Copy link


Copy Markdown

CollaboratorAuthor
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
Signed the commit. Thanks.
@ptrendx
ptrendx
commented
Nov 28, 2022

Copy link


Copy Markdown

Member
@ksivaman Could you review this? Thanks :-)!
vasunvidia
vasunvidia
commented
Dec 14, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

 CUBLASLT_MATMUL_DESC_AMAX_D_POINTER,
 &D_amax,
 sizeof(D_amax)));
 NVTE_CHECK_CUBLAS(cublasLtMatrixLayoutCreate(&Cdesc, bias_type, m, n, ldd));
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

What if C desc is same as D desc?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

For FP8 output, C cannot be FP8. So C type should be same as bias_type for FP8 output type and D_type for other output types.
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/pytorch/cpp_extensions.py
Comment threadtransformer_engine/pytorch/cpp_extensions.py
Outdated

 return_output = True

 out_dtype = tex.DType.kFloat32 if fp32_output else TE_DType[out_dtype]
 bias_dtype = output_dtype if bias is None else TE_DType[bias.dtype]
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

SHould not require bias_type in C api.
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Check with cublas team
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

@ksivaman Could you remind why bias_type should not be required?
Comment threadtransformer_engine/pytorch/cpp_extensions.py
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.h
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

You are leaking memory here since Cdesc was already created.
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Thanks for pointing it out. Will fix this.
ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

TBH I fail to see why do you have to set C descriptor - it is there for the beta=1 case, right? So it should always be of the same type as D?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

For FP8 GEMM output, cublas doesn't support FP8 C_type. So we need to set CDesc to bias type for FP8 output and D type for others. Does it make sense?
image
ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

This breaks beta=1 case that we use for wgrad accumulation, no?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Got it. Will address this.
vasunvidia
vasunvidia
commented
Jan 17, 2023
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

ptrendx
ptrendx
reviewed
Jan 31, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/ts_fp8_op.cpp
Outdated

@ptrendxptrendxJan 31, 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.

Why do you need to clone them?
@vasunvidiavasunvidiaJan 31, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

I just copied the usage for other scale factors such as

at::Tensor A_scale_inverse_arg = A_scale_inverse.clone();
Is it unnecessary?
@ptrendxptrendxJan 31, 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.

@ksivaman could you comment on that?
@ksivamanksivamanJan 31, 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.

Yes this shouldn't be cloned for 2 reasons:
  1. Unnecessary malloc,
  2. and more importantly, we want to populate the original amax tensor received in the arguments, which would remain unchanged here.
@vasunvidiavasunvidiaJan 31, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Thanks. I'll make the change.
@ptrendxptrendxFeb 1, 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.

@ksivaman That begs the question why are there the other clones for scale inverses? I believe those are unnecessary as well and create additional (albeit pretty small)D2D copies before every FP8 gemm.
@ksivamanksivamanFeb 1, 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.

Yes, I believe all 7 clones in transformer_engine/pytorch/csrc/ts_fp8_op.cpp can be removed. The scale inverse clones also seem unused? @asfiyab-nvidia Could you comment on why these were added to begin with?
@asfiyab-nvidiaasfiyab-nvidiaFeb 1, 2023

Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

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

@ksivaman These aren't necessary. I'll create a PR with a fix shortly. Thanks for pointing it out
ptrendx
ptrendx
reviewed
Jan 31, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/module.py
Outdated

@ptrendx
ptrendx
commented
Jan 31, 2023

Copy link


Copy Markdown

Member
/te-ci
@vasunvidia
Increase number of FP8 tensors per GEMM
95910d8
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Enable FP8 output tensor for fp8_gemm
fcbce36
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
[BERT FP8] Initial TE review comments
8a336ab
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Temporary fix for cuda graph non convergence
320182a
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Address review comments-2
ffe4852
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Review comments-3
075e287
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Cleanup
784fb47
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Change for New API
0ccb2e5
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Remove unnecessary clone for D_scale, D_amax
397a1e4
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Avoid Roll for AMAX history size = 1
a575b5e
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Update onnx_te_gemm API
2dad360
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
vasunvidia
commented
Feb 2, 2023

Copy link


Copy Markdown

CollaboratorAuthor
/te-ci

1 similar comment

@ptrendx
ptrendx
commented
Feb 2, 2023

Copy link


Copy Markdown

Member
/te-ci
@vasunvidia
Fix Lint errors
66b98e3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
ptrendx
ptrendx
approved these changes
Feb 3, 2023
View reviewed changes
@ptrendxptrendx
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, thanks :-).
@ptrendx
ptrendx
merged commit 14198f2
into
NVIDIA:mainFeb 3, 2023
cyanguwa pushed a commit
to cyanguwa/TransformerEngine
that referenced
this pull request
Feb 13, 2023
@vasunvidia@cyanguwa
Increase number of FP8 tensors per GEMM (NVIDIA#22)
f071530
* Increase number of FP8 tensors per GEMM
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Enable FP8 output tensor for fp8_gemm
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* [BERT FP8] Initial TE review comments
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Temporary fix for cuda graph non convergence
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Address review comments-2
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Review comments-3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Cleanup
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Change for New API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Remove unnecessary clone for D_scale, D_amax
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Avoid Roll for AMAX history size = 1
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Update onnx_te_gemm API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Fix Lint errors
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
---------
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
ksivaman added a commit
that referenced
this pull request
Feb 22, 2023
@cyanguwa@ksivaman
flash-attn integration (#62)
79a9fe2
* add flash attention to TransformerLayer
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Add docs for FP8 calibration (#61)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Fix the integer overflow in fused softmax (#60)
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* prefix flash attn env var with NVTE_
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Address steady memory increase and bloated checkpoints (#63)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix env var logic
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix flash attn env var logic again
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* remove d2d copies (#64)
* remove d2d copies
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* cleanup
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Increase number of FP8 tensors per GEMM (#22)
* Increase number of FP8 tensors per GEMM
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Enable FP8 output tensor for fp8_gemm
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* [BERT FP8] Initial TE review comments
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Temporary fix for cuda graph non convergence
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Address review comments-2
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Review comments-3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Cleanup
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Change for New API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Remove unnecessary clone for D_scale, D_amax
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Avoid Roll for AMAX history size = 1
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Update onnx_te_gemm API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Fix Lint errors
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
---------
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Bug fixes from PR 22 (#65)
* Bug fixes from PR 22
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add FP8 tests to ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* bundle unittests for ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* replace rearrange with transpose
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* QKV parameters unfused path fixes and optimization (#66)
* Bug fixes from PR 22
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add FP8 tests to ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Better QKV parameter fusion
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* small fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* keep original param for unfused case to retain externally set attrs
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Fix ONNX exports
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* improve arg naming
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* No need to set data pointers
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Assert memory loc in NoopCat
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Handle case of different memory in param and buffer
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix assert always true
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Reassign params memory to avoid more concats
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Fix gradients when using AMP (#70)
retain grad related attrs while casting
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix pylint violations fixed pyline violations such as trailing white spaces and too long lines Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* fix pylint violation on line 264 with R1719
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* fix two more pylint violations
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* DotProductAttention API
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add docs for attention
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix assert always true
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* check for correct flash-attn version
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* address review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint+build fixes, correct settings for default flash-attn
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* correct version
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review comments and fixes
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix onnx and disable flash-attn export test
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* remove einops dependency
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* cleanup internal API; rm duplication
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* only install TE wheel (exclude flash-attn to rm conflicts)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* forgot to change install wheel path
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* next round review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix flash_attn output
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix QK layer scaling
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* update docs
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review comments and fixes to selective checkpointing
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Charlene Yang <charleney@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Co-authored-by: Charlene Yang <charleney@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
shifangx pushed a commit
to shifangx/TransformerEngine
that referenced
this pull request
Feb 10, 2026
@lxd-cumt
Remove use_gems context and call flag_gems.xxx directly (NVIDIA#22)
832a797
- Remove the flag_gems.use_gems() context to avoid context-switching
overhead
- Call flag_gems.xxx directly wherever possible.
Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@ptrendxptrendxptrendx approved these changes
@ksivamanksivamanAwaiting requested review from ksivamanksivaman is a code owner

+1 more reviewer

@asfiyab-nvidiaasfiyab-nvidiaasfiyab-nvidia left review comments
Reviewers whose approvals may not affect merge requirements

Assignees


No one assigned


Labels


None yet


Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

4 participants

@vasunvidia@ptrendx@ksivaman@asfiyab-nvidia

, '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('^' + ".*" + '
Increase number of FP8 tensors per GEMM by vasunvidia · Pull Request #22 · NVIDIA/TransformerEngine · GitHub
Skip to content
Increase number of FP8 tensors per GEMM - #22
Merged
ptrendx merged 12 commits into
NVIDIA:mainfrom 
vasunvidia:main
Feb 3, 2023
Merged
Increase number of FP8 tensors per GEMM#22
ptrendx merged 12 commits into
NVIDIA:mainfrom 
vasunvidia:main
Conversation
@vasunvidia
@vasunvidiavasunvidia
commented
Oct 28, 2022

Copy link


Copy Markdown

Collaborator
 No description provided. 
@ptrendx
ptrendx
commented
Oct 28, 2022

Copy link


Copy Markdown

Member
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
@ptrendx
ptrendx
commented
Oct 28, 2022

Copy link


Copy Markdown

Member
/blossom-ci
@vasunvidia
vasunvidia
commented
Oct 29, 2022

Copy link


Copy Markdown

CollaboratorAuthor
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
Signed the commit. Thanks.
@ptrendx
ptrendx
commented
Nov 28, 2022

Copy link


Copy Markdown

Member
@ksivaman Could you review this? Thanks :-)!
vasunvidia
vasunvidia
commented
Dec 14, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

 CUBLASLT_MATMUL_DESC_AMAX_D_POINTER,
 &D_amax,
 sizeof(D_amax)));
 NVTE_CHECK_CUBLAS(cublasLtMatrixLayoutCreate(&Cdesc, bias_type, m, n, ldd));
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

What if C desc is same as D desc?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

For FP8 output, C cannot be FP8. So C type should be same as bias_type for FP8 output type and D_type for other output types.
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/pytorch/cpp_extensions.py
Comment threadtransformer_engine/pytorch/cpp_extensions.py
Outdated

 return_output = True

 out_dtype = tex.DType.kFloat32 if fp32_output else TE_DType[out_dtype]
 bias_dtype = output_dtype if bias is None else TE_DType[bias.dtype]
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

SHould not require bias_type in C api.
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Check with cublas team
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

@ksivaman Could you remind why bias_type should not be required?
Comment threadtransformer_engine/pytorch/cpp_extensions.py
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.h
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

You are leaking memory here since Cdesc was already created.
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Thanks for pointing it out. Will fix this.
ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

TBH I fail to see why do you have to set C descriptor - it is there for the beta=1 case, right? So it should always be of the same type as D?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

For FP8 GEMM output, cublas doesn't support FP8 C_type. So we need to set CDesc to bias type for FP8 output and D type for others. Does it make sense?
image
ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

This breaks beta=1 case that we use for wgrad accumulation, no?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Got it. Will address this.
vasunvidia
vasunvidia
commented
Jan 17, 2023
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

ptrendx
ptrendx
reviewed
Jan 31, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/ts_fp8_op.cpp
Outdated

@ptrendxptrendxJan 31, 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.

Why do you need to clone them?
@vasunvidiavasunvidiaJan 31, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

I just copied the usage for other scale factors such as

at::Tensor A_scale_inverse_arg = A_scale_inverse.clone();
Is it unnecessary?
@ptrendxptrendxJan 31, 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.

@ksivaman could you comment on that?
@ksivamanksivamanJan 31, 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.

Yes this shouldn't be cloned for 2 reasons:
  1. Unnecessary malloc,
  2. and more importantly, we want to populate the original amax tensor received in the arguments, which would remain unchanged here.
@vasunvidiavasunvidiaJan 31, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Thanks. I'll make the change.
@ptrendxptrendxFeb 1, 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.

@ksivaman That begs the question why are there the other clones for scale inverses? I believe those are unnecessary as well and create additional (albeit pretty small)D2D copies before every FP8 gemm.
@ksivamanksivamanFeb 1, 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.

Yes, I believe all 7 clones in transformer_engine/pytorch/csrc/ts_fp8_op.cpp can be removed. The scale inverse clones also seem unused? @asfiyab-nvidia Could you comment on why these were added to begin with?
@asfiyab-nvidiaasfiyab-nvidiaFeb 1, 2023

Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

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

@ksivaman These aren't necessary. I'll create a PR with a fix shortly. Thanks for pointing it out
ptrendx
ptrendx
reviewed
Jan 31, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/module.py
Outdated

@ptrendx
ptrendx
commented
Jan 31, 2023

Copy link


Copy Markdown

Member
/te-ci
@vasunvidia
Increase number of FP8 tensors per GEMM
95910d8
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Enable FP8 output tensor for fp8_gemm
fcbce36
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
[BERT FP8] Initial TE review comments
8a336ab
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Temporary fix for cuda graph non convergence
320182a
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Address review comments-2
ffe4852
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Review comments-3
075e287
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Cleanup
784fb47
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Change for New API
0ccb2e5
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Remove unnecessary clone for D_scale, D_amax
397a1e4
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Avoid Roll for AMAX history size = 1
a575b5e
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Update onnx_te_gemm API
2dad360
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
vasunvidia
commented
Feb 2, 2023

Copy link


Copy Markdown

CollaboratorAuthor
/te-ci

1 similar comment

@ptrendx
ptrendx
commented
Feb 2, 2023

Copy link


Copy Markdown

Member
/te-ci
@vasunvidia
Fix Lint errors
66b98e3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
ptrendx
ptrendx
approved these changes
Feb 3, 2023
View reviewed changes
@ptrendxptrendx
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, thanks :-).
@ptrendx
ptrendx
merged commit 14198f2
into
NVIDIA:mainFeb 3, 2023
cyanguwa pushed a commit
to cyanguwa/TransformerEngine
that referenced
this pull request
Feb 13, 2023
@vasunvidia@cyanguwa
Increase number of FP8 tensors per GEMM (NVIDIA#22)
f071530
* Increase number of FP8 tensors per GEMM
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Enable FP8 output tensor for fp8_gemm
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* [BERT FP8] Initial TE review comments
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Temporary fix for cuda graph non convergence
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Address review comments-2
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Review comments-3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Cleanup
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Change for New API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Remove unnecessary clone for D_scale, D_amax
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Avoid Roll for AMAX history size = 1
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Update onnx_te_gemm API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Fix Lint errors
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
---------
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
ksivaman added a commit
that referenced
this pull request
Feb 22, 2023
@cyanguwa@ksivaman
flash-attn integration (#62)
79a9fe2
* add flash attention to TransformerLayer
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Add docs for FP8 calibration (#61)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Fix the integer overflow in fused softmax (#60)
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* prefix flash attn env var with NVTE_
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Address steady memory increase and bloated checkpoints (#63)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix env var logic
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix flash attn env var logic again
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* remove d2d copies (#64)
* remove d2d copies
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* cleanup
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Increase number of FP8 tensors per GEMM (#22)
* Increase number of FP8 tensors per GEMM
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Enable FP8 output tensor for fp8_gemm
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* [BERT FP8] Initial TE review comments
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Temporary fix for cuda graph non convergence
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Address review comments-2
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Review comments-3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Cleanup
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Change for New API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Remove unnecessary clone for D_scale, D_amax
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Avoid Roll for AMAX history size = 1
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Update onnx_te_gemm API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Fix Lint errors
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
---------
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Bug fixes from PR 22 (#65)
* Bug fixes from PR 22
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add FP8 tests to ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* bundle unittests for ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* replace rearrange with transpose
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* QKV parameters unfused path fixes and optimization (#66)
* Bug fixes from PR 22
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add FP8 tests to ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Better QKV parameter fusion
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* small fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* keep original param for unfused case to retain externally set attrs
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Fix ONNX exports
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* improve arg naming
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* No need to set data pointers
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Assert memory loc in NoopCat
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Handle case of different memory in param and buffer
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix assert always true
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Reassign params memory to avoid more concats
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Fix gradients when using AMP (#70)
retain grad related attrs while casting
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix pylint violations fixed pyline violations such as trailing white spaces and too long lines Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* fix pylint violation on line 264 with R1719
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* fix two more pylint violations
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* DotProductAttention API
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add docs for attention
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix assert always true
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* check for correct flash-attn version
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* address review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint+build fixes, correct settings for default flash-attn
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* correct version
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review comments and fixes
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix onnx and disable flash-attn export test
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* remove einops dependency
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* cleanup internal API; rm duplication
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* only install TE wheel (exclude flash-attn to rm conflicts)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* forgot to change install wheel path
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* next round review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix flash_attn output
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix QK layer scaling
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* update docs
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review comments and fixes to selective checkpointing
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Charlene Yang <charleney@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Co-authored-by: Charlene Yang <charleney@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
shifangx pushed a commit
to shifangx/TransformerEngine
that referenced
this pull request
Feb 10, 2026
@lxd-cumt
Remove use_gems context and call flag_gems.xxx directly (NVIDIA#22)
832a797
- Remove the flag_gems.use_gems() context to avoid context-switching
overhead
- Call flag_gems.xxx directly wherever possible.
Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@ptrendxptrendxptrendx approved these changes
@ksivamanksivamanAwaiting requested review from ksivamanksivaman is a code owner

+1 more reviewer

@asfiyab-nvidiaasfiyab-nvidiaasfiyab-nvidia left review comments
Reviewers whose approvals may not affect merge requirements

Assignees


No one assigned


Labels


None yet


Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

4 participants

@vasunvidia@ptrendx@ksivaman@asfiyab-nvidia

, '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); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '
Increase number of FP8 tensors per GEMM by vasunvidia · Pull Request #22 · NVIDIA/TransformerEngine · GitHub
Skip to content
Increase number of FP8 tensors per GEMM - #22
Merged
ptrendx merged 12 commits into
NVIDIA:mainfrom 
vasunvidia:main
Feb 3, 2023
Merged
Increase number of FP8 tensors per GEMM#22
ptrendx merged 12 commits into
NVIDIA:mainfrom 
vasunvidia:main
Conversation
@vasunvidia
@vasunvidiavasunvidia
commented
Oct 28, 2022

Copy link


Copy Markdown

Collaborator
 No description provided. 
@ptrendx
ptrendx
commented
Oct 28, 2022

Copy link


Copy Markdown

Member
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
@ptrendx
ptrendx
commented
Oct 28, 2022

Copy link


Copy Markdown

Member
/blossom-ci
@vasunvidia
vasunvidia
commented
Oct 29, 2022

Copy link


Copy Markdown

CollaboratorAuthor
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
Signed the commit. Thanks.
@ptrendx
ptrendx
commented
Nov 28, 2022

Copy link


Copy Markdown

Member
@ksivaman Could you review this? Thanks :-)!
vasunvidia
vasunvidia
commented
Dec 14, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

 CUBLASLT_MATMUL_DESC_AMAX_D_POINTER,
 &D_amax,
 sizeof(D_amax)));
 NVTE_CHECK_CUBLAS(cublasLtMatrixLayoutCreate(&Cdesc, bias_type, m, n, ldd));
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

What if C desc is same as D desc?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

For FP8 output, C cannot be FP8. So C type should be same as bias_type for FP8 output type and D_type for other output types.
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/pytorch/cpp_extensions.py
Comment threadtransformer_engine/pytorch/cpp_extensions.py
Outdated

 return_output = True

 out_dtype = tex.DType.kFloat32 if fp32_output else TE_DType[out_dtype]
 bias_dtype = output_dtype if bias is None else TE_DType[bias.dtype]
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

SHould not require bias_type in C api.
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Check with cublas team
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

@ksivaman Could you remind why bias_type should not be required?
Comment threadtransformer_engine/pytorch/cpp_extensions.py
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.h
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

You are leaking memory here since Cdesc was already created.
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Thanks for pointing it out. Will fix this.
ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

TBH I fail to see why do you have to set C descriptor - it is there for the beta=1 case, right? So it should always be of the same type as D?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

For FP8 GEMM output, cublas doesn't support FP8 C_type. So we need to set CDesc to bias type for FP8 output and D type for others. Does it make sense?
image
ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

This breaks beta=1 case that we use for wgrad accumulation, no?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Got it. Will address this.
vasunvidia
vasunvidia
commented
Jan 17, 2023
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

ptrendx
ptrendx
reviewed
Jan 31, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/ts_fp8_op.cpp
Outdated

@ptrendxptrendxJan 31, 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.

Why do you need to clone them?
@vasunvidiavasunvidiaJan 31, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

I just copied the usage for other scale factors such as

at::Tensor A_scale_inverse_arg = A_scale_inverse.clone();
Is it unnecessary?
@ptrendxptrendxJan 31, 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.

@ksivaman could you comment on that?
@ksivamanksivamanJan 31, 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.

Yes this shouldn't be cloned for 2 reasons:
  1. Unnecessary malloc,
  2. and more importantly, we want to populate the original amax tensor received in the arguments, which would remain unchanged here.
@vasunvidiavasunvidiaJan 31, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Thanks. I'll make the change.
@ptrendxptrendxFeb 1, 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.

@ksivaman That begs the question why are there the other clones for scale inverses? I believe those are unnecessary as well and create additional (albeit pretty small)D2D copies before every FP8 gemm.
@ksivamanksivamanFeb 1, 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.

Yes, I believe all 7 clones in transformer_engine/pytorch/csrc/ts_fp8_op.cpp can be removed. The scale inverse clones also seem unused? @asfiyab-nvidia Could you comment on why these were added to begin with?
@asfiyab-nvidiaasfiyab-nvidiaFeb 1, 2023

Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

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

@ksivaman These aren't necessary. I'll create a PR with a fix shortly. Thanks for pointing it out
ptrendx
ptrendx
reviewed
Jan 31, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/module.py
Outdated

@ptrendx
ptrendx
commented
Jan 31, 2023

Copy link


Copy Markdown

Member
/te-ci
@vasunvidia
Increase number of FP8 tensors per GEMM
95910d8
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Enable FP8 output tensor for fp8_gemm
fcbce36
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
[BERT FP8] Initial TE review comments
8a336ab
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Temporary fix for cuda graph non convergence
320182a
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Address review comments-2
ffe4852
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Review comments-3
075e287
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Cleanup
784fb47
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Change for New API
0ccb2e5
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Remove unnecessary clone for D_scale, D_amax
397a1e4
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Avoid Roll for AMAX history size = 1
a575b5e
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Update onnx_te_gemm API
2dad360
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
vasunvidia
commented
Feb 2, 2023

Copy link


Copy Markdown

CollaboratorAuthor
/te-ci

1 similar comment

@ptrendx
ptrendx
commented
Feb 2, 2023

Copy link


Copy Markdown

Member
/te-ci
@vasunvidia
Fix Lint errors
66b98e3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
ptrendx
ptrendx
approved these changes
Feb 3, 2023
View reviewed changes
@ptrendxptrendx
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, thanks :-).
@ptrendx
ptrendx
merged commit 14198f2
into
NVIDIA:mainFeb 3, 2023
cyanguwa pushed a commit
to cyanguwa/TransformerEngine
that referenced
this pull request
Feb 13, 2023
@vasunvidia@cyanguwa
Increase number of FP8 tensors per GEMM (NVIDIA#22)
f071530
* Increase number of FP8 tensors per GEMM
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Enable FP8 output tensor for fp8_gemm
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* [BERT FP8] Initial TE review comments
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Temporary fix for cuda graph non convergence
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Address review comments-2
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Review comments-3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Cleanup
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Change for New API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Remove unnecessary clone for D_scale, D_amax
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Avoid Roll for AMAX history size = 1
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Update onnx_te_gemm API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Fix Lint errors
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
---------
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
ksivaman added a commit
that referenced
this pull request
Feb 22, 2023
@cyanguwa@ksivaman
flash-attn integration (#62)
79a9fe2
* add flash attention to TransformerLayer
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Add docs for FP8 calibration (#61)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Fix the integer overflow in fused softmax (#60)
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* prefix flash attn env var with NVTE_
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Address steady memory increase and bloated checkpoints (#63)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix env var logic
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix flash attn env var logic again
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* remove d2d copies (#64)
* remove d2d copies
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* cleanup
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Increase number of FP8 tensors per GEMM (#22)
* Increase number of FP8 tensors per GEMM
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Enable FP8 output tensor for fp8_gemm
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* [BERT FP8] Initial TE review comments
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Temporary fix for cuda graph non convergence
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Address review comments-2
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Review comments-3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Cleanup
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Change for New API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Remove unnecessary clone for D_scale, D_amax
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Avoid Roll for AMAX history size = 1
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Update onnx_te_gemm API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Fix Lint errors
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
---------
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Bug fixes from PR 22 (#65)
* Bug fixes from PR 22
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add FP8 tests to ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* bundle unittests for ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* replace rearrange with transpose
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* QKV parameters unfused path fixes and optimization (#66)
* Bug fixes from PR 22
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add FP8 tests to ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Better QKV parameter fusion
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* small fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* keep original param for unfused case to retain externally set attrs
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Fix ONNX exports
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* improve arg naming
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* No need to set data pointers
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Assert memory loc in NoopCat
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Handle case of different memory in param and buffer
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix assert always true
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Reassign params memory to avoid more concats
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Fix gradients when using AMP (#70)
retain grad related attrs while casting
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix pylint violations fixed pyline violations such as trailing white spaces and too long lines Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* fix pylint violation on line 264 with R1719
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* fix two more pylint violations
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* DotProductAttention API
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add docs for attention
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix assert always true
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* check for correct flash-attn version
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* address review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint+build fixes, correct settings for default flash-attn
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* correct version
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review comments and fixes
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix onnx and disable flash-attn export test
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* remove einops dependency
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* cleanup internal API; rm duplication
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* only install TE wheel (exclude flash-attn to rm conflicts)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* forgot to change install wheel path
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* next round review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix flash_attn output
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix QK layer scaling
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* update docs
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review comments and fixes to selective checkpointing
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Charlene Yang <charleney@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Co-authored-by: Charlene Yang <charleney@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
shifangx pushed a commit
to shifangx/TransformerEngine
that referenced
this pull request
Feb 10, 2026
@lxd-cumt
Remove use_gems context and call flag_gems.xxx directly (NVIDIA#22)
832a797
- Remove the flag_gems.use_gems() context to avoid context-switching
overhead
- Call flag_gems.xxx directly wherever possible.
Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@ptrendxptrendxptrendx approved these changes
@ksivamanksivamanAwaiting requested review from ksivamanksivaman is a code owner

+1 more reviewer

@asfiyab-nvidiaasfiyab-nvidiaasfiyab-nvidia left review comments
Reviewers whose approvals may not affect merge requirements

Assignees


No one assigned


Labels


None yet


Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

4 participants

@vasunvidia@ptrendx@ksivaman@asfiyab-nvidia

, 'i');
if (__m === '*' || __re.test(location.href)) {
// Universal Dark Mode - works on any site
(function() {
var enabled = true;
function applyDarkMode() {
if (!enabled) return;
// Create style element if it doesn't exist
var style = document.getElementById('universal-dark-mode-style');
if (!style) {
style = document.createElement('style');
style.id = 'universal-dark-mode-style';
document.head.appendChild(style);
}
// Dark mode CSS - inverts colors but preserves images/video
style.textContent = '
/* Invert everything except media */
html {
filter: invert(1) hue-rotate(180deg) !important;
background: #1a1a2e !important;
}
/* Restore images, videos, iframes, canvas */
img, video, iframe, canvas, svg, picture, [style*="background-image"] {
filter: invert(1) hue-rotate(180deg) !important;
}
/* Preserve specific elements that should not be inverted */
.no-dark-mode, .no-dark-mode *,
[data-theme="light"], [data-theme="light"],
.ace_editor, .ace_editor *,
.CodeMirror, .CodeMirror *,
.monaco-editor, .monaco-editor *,
.markdown-body pre, .markdown-body pre *,
.highlight, .highlight *,
pre code, pre code * {
filter: none !important;
}
/* Fix common UI elements */
.modal, .popup, .dropdown-menu, .tooltip, .popover {
filter: invert(1) hue-rotate(180deg) !important;
background: #2d2d44 !important;
border-color: #444 !important;
}
/* Scrollbars */
::-webkit-scrollbar { background: #1a1a2e !important; }
::-webkit-scrollbar-thumb { background: #444 !important; }
::-webkit-scrollbar-thumb:hover { background: #555 !important; }
/* Selection */
::selection { background: #4ecdc4 !important; color: #1a1a2e !important; }
::-moz-selection { background: #4ecdc4 !important; color: #1a1a2e !important; }
';
}
function removeDarkMode() {
var style = document.getElementById('universal-dark-mode-style');
if (style) style.remove();
}
// Toggle with Alt+Shift+D
document.addEventListener('keydown', function(e) {
if (e.altKey && e.shiftKey && e.key === 'D') {
e.preventDefault();
enabled = !enabled;
if (enabled) {
applyDarkMode();
console.log('[Universal Dark Mode] Enabled');
} else {
removeDarkMode();
console.log('[Universal Dark Mode] Disabled');
}
}
});
// Apply on load
applyDarkMode();
// Re-apply on dynamic content
var observer = new MutationObserver(function(mutations) {
if (enabled && !document.getElementById('universal-dark-mode-style')) {
applyDarkMode();
}
});
observer.observe(document.head, { childList: true });
console.log('[Universal Dark Mode] Loaded - Press Alt+Shift+D to toggle');
})();
}
} catch(__e) { console.warn('[Userscript:Universal Dark Mode]', __e); }
})();
})();
Increase number of FP8 tensors per GEMM by vasunvidia · Pull Request #22 · NVIDIA/TransformerEngine · GitHub
Skip to content
Increase number of FP8 tensors per GEMM - #22
Merged
ptrendx merged 12 commits into
NVIDIA:mainfrom 
vasunvidia:main
Feb 3, 2023
Merged
Increase number of FP8 tensors per GEMM#22
ptrendx merged 12 commits into
NVIDIA:mainfrom 
vasunvidia:main
Conversation
@vasunvidia
@vasunvidiavasunvidia
commented
Oct 28, 2022

Copy link


Copy Markdown

Collaborator
 No description provided. 
@ptrendx
ptrendx
commented
Oct 28, 2022

Copy link


Copy Markdown

Member
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
@ptrendx
ptrendx
commented
Oct 28, 2022

Copy link


Copy Markdown

Member
/blossom-ci
@vasunvidia
vasunvidia
commented
Oct 29, 2022

Copy link


Copy Markdown

CollaboratorAuthor
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
@vasunvidia Please sign your commits (see CONTRIBUTING.rst)
Signed the commit. Thanks.
@ptrendx
ptrendx
commented
Nov 28, 2022

Copy link


Copy Markdown

Member
@ksivaman Could you review this? Thanks :-)!
vasunvidia
vasunvidia
commented
Dec 14, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

 CUBLASLT_MATMUL_DESC_AMAX_D_POINTER,
 &D_amax,
 sizeof(D_amax)));
 NVTE_CHECK_CUBLAS(cublasLtMatrixLayoutCreate(&Cdesc, bias_type, m, n, ldd));
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

What if C desc is same as D desc?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

For FP8 output, C cannot be FP8. So C type should be same as bias_type for FP8 output type and D_type for other output types.
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

Comment threadtransformer_engine/pytorch/cpp_extensions.py
Comment threadtransformer_engine/pytorch/cpp_extensions.py
Outdated

 return_output = True

 out_dtype = tex.DType.kFloat32 if fp32_output else TE_DType[out_dtype]
 bias_dtype = output_dtype if bias is None else TE_DType[bias.dtype]
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

SHould not require bias_type in C api.
@vasunvidiavasunvidiaDec 14, 2022

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Check with cublas team
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

@ksivaman Could you remind why bias_type should not be required?
Comment threadtransformer_engine/pytorch/cpp_extensions.py
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.h
Outdated

Comment threadtransformer_engine/pytorch/csrc/extensions.cu
Outdated

ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

You are leaking memory here since Cdesc was already created.
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Thanks for pointing it out. Will fix this.
ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

TBH I fail to see why do you have to set C descriptor - it is there for the beta=1 case, right? So it should always be of the same type as D?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

For FP8 GEMM output, cublas doesn't support FP8 C_type. So we need to set CDesc to bias type for FP8 output and D type for others. Does it make sense?
image
ptrendx
ptrendx
reviewed
Dec 16, 2022
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

@ptrendxptrendxDec 16, 2022

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.

This breaks beta=1 case that we use for wgrad accumulation, no?
@vasunvidiavasunvidiaJan 8, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Got it. Will address this.
vasunvidia
vasunvidia
commented
Jan 17, 2023
View reviewed changes
Comment threadtransformer_engine/common/gemm/cublaslt_gemm.cu
Outdated

ptrendx
ptrendx
reviewed
Jan 31, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/csrc/ts_fp8_op.cpp
Outdated

@ptrendxptrendxJan 31, 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.

Why do you need to clone them?
@vasunvidiavasunvidiaJan 31, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

I just copied the usage for other scale factors such as

at::Tensor A_scale_inverse_arg = A_scale_inverse.clone();
Is it unnecessary?
@ptrendxptrendxJan 31, 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.

@ksivaman could you comment on that?
@ksivamanksivamanJan 31, 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.

Yes this shouldn't be cloned for 2 reasons:
  1. Unnecessary malloc,
  2. and more importantly, we want to populate the original amax tensor received in the arguments, which would remain unchanged here.
@vasunvidiavasunvidiaJan 31, 2023

Copy link


Copy Markdown

CollaboratorAuthor
Choose a reason for hiding this comment

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

Thanks. I'll make the change.
@ptrendxptrendxFeb 1, 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.

@ksivaman That begs the question why are there the other clones for scale inverses? I believe those are unnecessary as well and create additional (albeit pretty small)D2D copies before every FP8 gemm.
@ksivamanksivamanFeb 1, 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.

Yes, I believe all 7 clones in transformer_engine/pytorch/csrc/ts_fp8_op.cpp can be removed. The scale inverse clones also seem unused? @asfiyab-nvidia Could you comment on why these were added to begin with?
@asfiyab-nvidiaasfiyab-nvidiaFeb 1, 2023

Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

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

@ksivaman These aren't necessary. I'll create a PR with a fix shortly. Thanks for pointing it out
ptrendx
ptrendx
reviewed
Jan 31, 2023
View reviewed changes
Comment threadtransformer_engine/pytorch/module.py
Outdated

@ptrendx
ptrendx
commented
Jan 31, 2023

Copy link


Copy Markdown

Member
/te-ci
@vasunvidia
Increase number of FP8 tensors per GEMM
95910d8
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Enable FP8 output tensor for fp8_gemm
fcbce36
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
[BERT FP8] Initial TE review comments
8a336ab
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Temporary fix for cuda graph non convergence
320182a
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Address review comments-2
ffe4852
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Review comments-3
075e287
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Cleanup
784fb47
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Change for New API
0ccb2e5
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Remove unnecessary clone for D_scale, D_amax
397a1e4
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Avoid Roll for AMAX history size = 1
a575b5e
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
Update onnx_te_gemm API
2dad360
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
@vasunvidia
vasunvidia
commented
Feb 2, 2023

Copy link


Copy Markdown

CollaboratorAuthor
/te-ci

1 similar comment

@ptrendx
ptrendx
commented
Feb 2, 2023

Copy link


Copy Markdown

Member
/te-ci
@vasunvidia
Fix Lint errors
66b98e3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
ptrendx
ptrendx
approved these changes
Feb 3, 2023
View reviewed changes
@ptrendxptrendx
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, thanks :-).
@ptrendx
ptrendx
merged commit 14198f2
into
NVIDIA:mainFeb 3, 2023
cyanguwa pushed a commit
to cyanguwa/TransformerEngine
that referenced
this pull request
Feb 13, 2023
@vasunvidia@cyanguwa
Increase number of FP8 tensors per GEMM (NVIDIA#22)
f071530
* Increase number of FP8 tensors per GEMM
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Enable FP8 output tensor for fp8_gemm
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* [BERT FP8] Initial TE review comments
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Temporary fix for cuda graph non convergence
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Address review comments-2
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Review comments-3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Cleanup
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Change for New API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Remove unnecessary clone for D_scale, D_amax
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Avoid Roll for AMAX history size = 1
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Update onnx_te_gemm API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Fix Lint errors
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
---------
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
ksivaman added a commit
that referenced
this pull request
Feb 22, 2023
@cyanguwa@ksivaman
flash-attn integration (#62)
79a9fe2
* add flash attention to TransformerLayer
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Add docs for FP8 calibration (#61)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Fix the integer overflow in fused softmax (#60)
Signed-off-by: Przemek Tredak <ptredak@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* prefix flash attn env var with NVTE_
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Address steady memory increase and bloated checkpoints (#63)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix env var logic
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix flash attn env var logic again
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* remove d2d copies (#64)
* remove d2d copies
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* cleanup
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Increase number of FP8 tensors per GEMM (#22)
* Increase number of FP8 tensors per GEMM
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Enable FP8 output tensor for fp8_gemm
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* [BERT FP8] Initial TE review comments
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Temporary fix for cuda graph non convergence
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Address review comments-2
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Review comments-3
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Cleanup
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Change for New API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Remove unnecessary clone for D_scale, D_amax
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Avoid Roll for AMAX history size = 1
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Update onnx_te_gemm API
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
* Fix Lint errors
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
---------
Signed-off-by: Vasudevan Rengasamy <vrengasamy@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Bug fixes from PR 22 (#65)
* Bug fixes from PR 22
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add FP8 tests to ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* bundle unittests for ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* replace rearrange with transpose
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* QKV parameters unfused path fixes and optimization (#66)
* Bug fixes from PR 22
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add FP8 tests to ci
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Better QKV parameter fusion
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* small fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* keep original param for unfused case to retain externally set attrs
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Fix ONNX exports
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* improve arg naming
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* No need to set data pointers
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Assert memory loc in NoopCat
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Handle case of different memory in param and buffer
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix assert always true
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Reassign params memory to avoid more concats
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* Fix gradients when using AMP (#70)
retain grad related attrs while casting
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: Charlene Yang <charleney@nvidia.com>
* fix pylint violations fixed pyline violations such as trailing white spaces and too long lines Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* fix pylint violation on line 264 with R1719
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* fix two more pylint violations
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
* DotProductAttention API
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* Add docs for attention
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix assert always true
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* check for correct flash-attn version
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* address review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* lint+build fixes, correct settings for default flash-attn
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* correct version
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review comments and fixes
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix onnx and disable flash-attn export test
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* remove einops dependency
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* cleanup internal API; rm duplication
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* only install TE wheel (exclude flash-attn to rm conflicts)
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* forgot to change install wheel path
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* next round review comments
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix flash_attn output
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* fix QK layer scaling
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* update docs
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
* review comments and fixes to selective checkpointing
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
---------
Signed-off-by: Charlene Yang <charleney@nvidia.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Signed-off-by: cyanguwa <cyang.uwa@gmail.com>
Co-authored-by: Charlene Yang <charleney@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
shifangx pushed a commit
to shifangx/TransformerEngine
that referenced
this pull request
Feb 10, 2026
@lxd-cumt
Remove use_gems context and call flag_gems.xxx directly (NVIDIA#22)
832a797
- Remove the flag_gems.use_gems() context to avoid context-switching
overhead
- Call flag_gems.xxx directly wherever possible.
Sign up for freeto join this conversation on GitHub.
Already have an account?
Sign in to comment

Reviewers

@ptrendxptrendxptrendx approved these changes
@ksivamanksivamanAwaiting requested review from ksivamanksivaman is a code owner

+1 more reviewer

@asfiyab-nvidiaasfiyab-nvidiaasfiyab-nvidia left review comments
Reviewers whose approvals may not affect merge requirements

Assignees


No one assigned


Labels


None yet


Projects


None yet


Milestone


No milestone


Development

Successfully merging this pull request may close these issues.

4 participants

@vasunvidia@ptrendx@ksivaman@asfiyab-nvidia