Commits · 20065c44a00f4a3bfce621ec59b7ba593b35976c · OpenDAS / TransformerEngine

21 Aug, 2025 1 commit
- [DCU] tensorwise int8 train opt · 7a923605
  yuguo authored Aug 21, 2025
  
  7a923605
13 Aug, 2025 2 commits
- Formatted · c6dae0e5
  wenjh authored Aug 13, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  c6dae0e5
- Delete tmpArgs in groupedgemm · 3cdceb87
  wenjh authored Aug 13, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  3cdceb87
12 Aug, 2025 1 commit
- Use hipMemcpy rather than hipMemcpyAsync · 1f97aebb
  wenjh authored Aug 12, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  1f97aebb
11 Aug, 2025 1 commit
- Fix hipblaslt_goupedgemm · 8665c111
  wenjh authored Aug 11, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  8665c111
08 Aug, 2025 2 commits
- [DCU] add tmp_userArgs for blaslt groupgemm · 14aaa48a
  yuguo authored Aug 08, 2025
  
  14aaa48a
- [DCU] add NVTE_INT8_SIM_FP8_TENSORWISE · 3eb6ea62
  yuguo authored Aug 08, 2025
  
  3eb6ea62
06 Aug, 2025 3 commits
- [DCU] fix channelwise train oom bug · 4a013bd5
  yuguo authored Aug 06, 2025
  
  4a013bd5
- [DCU] userargs groupedgemm · ddfbdaf4
  yuguo authored Aug 06, 2025
  
  ddfbdaf4
- [DCU] fix batchlinear core dump in 2.5 · eac0d49b
  yuguo authored Aug 06, 2025
  
  eac0d49b
22 Jul, 2025 2 commits
- [DCU] fix compile · 782f6092
  yuguo authored Jul 22, 2025
  
  782f6092
- [DCU] fix merge · eac75188
  yuguo authored Jul 22, 2025
  
  eac75188
18 Jul, 2025 2 commits

[DCU] fix bias gradient · 736e8f8b
yuguo authored Jul 18, 2025

736e8f8b

Update cudnn-frontend to 1.13.0 (#1960) · 657c965b

Charlene Yang authored Jul 17, 2025



* update cudnn-frontend to 1.13.0
Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* disable 9.11 for a bug
Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci



* fix selection logic
Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci



---------
Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>

657c965b

17 Jul, 2025 2 commits

[Common] Optimize KV cache related kernels (#1914) · f8933bbf

Charlene Yang authored Jul 17, 2025



* optimize kv_cache reindex and copy kernels
Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci



* avoid reindexing from python side
Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* rename variable from previous commit
Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* minor fix
Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

* minor fix
Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>

---------
Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>

f8933bbf

[DCU] fix · 00738a42
yuguo authored Jul 17, 2025

00738a42

16 Jul, 2025 2 commits

Handle dtypes more carefully in multi-tensor Adam (#1888) · bda29934

Tim Moon authored Jul 16, 2025



* Add dtype checks in multi-tensor Adam
Signed-off-by: Tim Moon <tmoon@nvidia.com>

* Avoid throwing exception in destructor
Signed-off-by: Tim Moon <tmoon@nvidia.com>

* Debug test failures
Signed-off-by: Tim Moon <tmoon@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci



---------
Signed-off-by: Tim Moon <tmoon@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>

bda29934

[DCU] surpport NVTE_USE_HIPBLASLT_GROUPEDGEMM · 9406ff31
yuguo authored Jul 16, 2025

9406ff31

15 Jul, 2025 1 commit
- [DCU] surpport channelwise int8 train · 793e0103
  yuguo authored Jul 15, 2025
  
  793e0103
14 Jul, 2025 2 commits

Run-time checks for CUDA and cuBLAS versions (#1938) · 1c702b4c

Tim Moon authored Jul 14, 2025



* Add run-time version checks in cuBLAS GEMM wrapper
Signed-off-by: Tim Moon <tmoon@nvidia.com>

* Add run-time version logic for multicast
Signed-off-by: Tim Moon <tmoon@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci



* Fix namespace error
Signed-off-by: Tim Moon <tmoon@nvidia.com>

---------
Signed-off-by: Tim Moon <tmoon@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>

1c702b4c

[PyTorch] Fix bugs in router fusion (#1944) · 397c4be6

Autumn1998 authored Jul 14, 2025



* fix underterminsic problem in CI
Signed-off-by: tongliu <tongliu@nvidia.com>

* fix bug on mbs>1
Signed-off-by: tongliu <tongliu@nvidia.com>

* fix bug on sm dispatcher
Signed-off-by: tongliu <tongliu@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci



* fix CI initial values
Signed-off-by: tongliu <tongliu@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci



---------
Signed-off-by: tongliu <tongliu@nvidia.com>
Co-authored-by: tongliu <tongliu@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Co-authored-by: Xin Yao <xiny@nvidia.com>

397c4be6

12 Jul, 2025 1 commit

Add backward fusions of dbias+quantize and dbias+dactivation+quantize to `te.Sequential` (#1942) · 37da2d3b

Jan Bielak authored Jul 11, 2025



* Fix clearing tensor data in backward removing is_first_op
Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* Misc fixes
Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* Use Linear weight dtype and device for compute consistently
Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* Add backward dbias + quantize fusion
Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* Pass recipe to OperationFuser to allow recipe-dependent fusions
Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* Remove redundant view from activations
Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* Add bias activation backward fusion
Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* Apply suggestions from code review
Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci



---------
Signed-off-by: Jan Bielak <jbielak@nvidia.com>
Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>

37da2d3b

10 Jul, 2025 1 commit

[PyTorch][MoE] Kernels fusions for the MoE router (#1883) · 62acae05

Autumn1998 authored Jul 11, 2025



* add router fusion
Signed-off-by: tongliu <tongliu@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci



* fix ci
Signed-off-by: tongliu <tongliu@nvidia.com>

* fix ci with cuda 12.3
Signed-off-by: tongliu <tongliu@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci



* Review suggestions
Signed-off-by: Tim Moon <tmoon@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci



* Fix CI sm89/80
Signed-off-by: tongliu <tongliu@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci



---------
Signed-off-by: tongliu <tongliu@nvidia.com>
Signed-off-by: Tim Moon <tmoon@nvidia.com>
Co-authored-by: tongliu <tongliu@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Co-authored-by: Xin Yao <xiny@nvidia.com>
Co-authored-by: Tim Moon <tmoon@nvidia.com>

62acae05

09 Jul, 2025 1 commit
- [DCU] channelwise batchgemm for MOE · 76023d21
  yuguo authored Jul 09, 2025
  
  76023d21
08 Jul, 2025 1 commit
- [DCU] Preliminary support for channelwise · 9fe13a33
  yuguo authored Jul 08, 2025
  
  9fe13a33
03 Jul, 2025 1 commit
- Fix kernel crash on block_len=64 · 40a4d896
  wenjh authored Jul 03, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  40a4d896
02 Jul, 2025 1 commit
- Resolve merge issues from develop_v2.4 · 4ef4eae6
  wenjh authored Jul 02, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  4ef4eae6
01 Jul, 2025 1 commit

[Blockwise] Add support block_len=64 support · b944277c

wenjh authored Jun 25, 2025



Add env to chose blocklen of blockwise quantize.
Signed-off-by: wenjh <wenjh@sugon.com>

Fix pytest of blockwise error
Signed-off-by: wenjh <wenjh@sugon.com>

Resolve new api in  int8 gemm test
Signed-off-by: wenjh <wenjh@sugon.com>

Fix incorrect launch parm
Signed-off-by: wenjh <wenjh@sugon.com>

Fix 1D blockwise(64) acc error
Signed-off-by: wenjh <wenjh@sugon.com>

b944277c

26 Jun, 2025 1 commit

[PyTorch|common] Optimize unpadding kernel for FP8 (#1866) · 23cf4ff9

xiaoxi-wangfj authored Jun 26, 2025



* [PyTorch|common] Implement unpadding kernel for FP8

1. Add multi-tensor unpadding kernel
2. Replace split+cat with unpadding kernel in Fp8Padding and Fp8Unpadding
3. Add unpadding with padding unit tests
Signed-off-by: xiaoxi-wangfj <690912414@qq.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci



* add license
Signed-off-by: Xin Yao <xiny@nvidia.com>

* Update padding.cu
Signed-off-by: Xin Yao <xiny@nvidia.com>

---------
Signed-off-by: xiaoxi-wangfj <690912414@qq.com>
Signed-off-by: Xin Yao <xiny@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Co-authored-by: Xin Yao <xiny@nvidia.com>

23cf4ff9

20 Jun, 2025 1 commit
- [DCU] fix 2.5 compile issues · 68487b2a
  yuguo authored Jun 20, 2025
  
  68487b2a
19 Jun, 2025 2 commits
- [DCU] fix TORCH_COMM_CU_NUMS conflicts in 2.5 · ffbef335
  yuguo authored Jun 19, 2025
  
  ffbef335
- [DCU] add TORCH_COMM_CU_NUMS and fix · d6c32078
  yuguo authored Jun 19, 2025
  
  d6c32078
18 Jun, 2025 5 commits
- Fix build error · e704bbc8
  wenjh authored Jun 18, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  e704bbc8
- Resolve merge issue from nv of vector blockwise · 331f2fc4
  wenjh authored Jun 18, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  331f2fc4
- Fix vector blockwise acc problem · 8a03ff34
  wenjh authored Jun 18, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  8a03ff34
- [DCU] fix 2.5 compile issues · 2b1428ff
  yuguo authored Jun 18, 2025
  
  2b1428ff
- Fix lack of lds in vector_blockwise · d1bf39cf
  wenjh authored Jun 18, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  d1bf39cf
17 Jun, 2025 1 commit
- [DCU] fix 2.5 get_num_compute_streams · b4a2489f
  yuguo authored Jun 17, 2025
  
  b4a2489f
16 Jun, 2025 1 commit

[JAX] Grouped GEMM & Dense support MXFP8 and handle empty matrices (#1871) · 01a504c4

Hua Huang authored Jun 16, 2025



* Support MXFP8 and handle empty matrices
Signed-off-by: Hua Huang <huah@nvidia.com>

---------
Signed-off-by: Hua Huang <huah@nvidia.com>

01a504c4

13 Jun, 2025 1 commit

Add support for head_dim > 128 (#1797) · 71c76b6b

Charlene Yang authored Jun 14, 2025



* add support for head dim > 128
Signed-off-by: Charlene Yang <charleney@nvidia.com>

* remove debugging
Signed-off-by: Charlene Yang <charleney@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci



* raise tols slightly to tolerate 1/2048 mismatches
Signed-off-by: Charlene Yang <charleney@nvidia.com>

* fix is_training for test_te_layer
Signed-off-by: Charlene Yang <charleney@nvidia.com>

* add bprop support for blackwell
Signed-off-by: Charlene Yang <charleney@nvidia.com>

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci



* minor tweak for format
Signed-off-by: Charlene Yang <charleney@nvidia.com>

* fix backend selection results
Signed-off-by: Charlene Yang <charleney@nvidia.com>

* bump sm100 to sm100+
Signed-off-by: Charlene Yang <charleney@nvidia.com>

* add sq=1 test for MLA
Signed-off-by: Charlene Yang <charleney@nvidia.com>

* enable sq=1 for bprop
Signed-off-by: Charlene Yang <charleney@nvidia.com>

* minor tweak in comments
Signed-off-by: Charlene Yang <charleney@nvidia.com>

* fix head_dim logic and remove pytest skip
Signed-off-by: Charlene Yang <charleney@nvidia.com>

* add FE fix for d>128
Signed-off-by: Charlene Yang <charleney@nvidia.com>

* update FE again to take in small fixes
Signed-off-by: Charlene Yang <charleney@nvidia.com>

* add cuDNN version info in L0 tests
Signed-off-by: Charlene Yang <charleney@nvidia.com>

* increase tols for Unfused + large dim
Signed-off-by: Charlene Yang <charleney@nvidia.com>

* Revert "add cuDNN version info in L0 tests"

This reverts commit 3e1b426ca5319a2c0540b9e73bba7047d0e583e5.
Signed-off-by: Charlene Yang <charleney@nvidia.com>

* fix tols for Unfused
Signed-off-by: Charlene Yang <charleney@nvidia.com>

---------
Signed-off-by: Charlene Yang <charleney@nvidia.com>
Signed-off-by: Charlene Yang <8636796+cyanguwa@users.noreply.github.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>

71c76b6b