Commits · e908bd30397423e3713318c63e82612be0097b06 · OpenDAS / TransformerEngine

07 Jan, 2026 1 commit
- Add nmz support · e908bd30
  wenjh authored Jan 07, 2026
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  e908bd30
19 Dec, 2025 1 commit
- Fix build error · a9601800
  wenjh authored Dec 19, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  a9601800
18 Dec, 2025 1 commit
- Add bias fwd/bwd at group gemm · 5cf21c3b
  wenjh authored Dec 18, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  5cf21c3b
15 Dec, 2025 1 commit
- Complete manage blaslt handle · 86d5cd03
  wenjh authored Dec 15, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  86d5cd03
13 Dec, 2025 1 commit
- Complete fix blaslt group gemm dump · 1261da47
  wenjh authored Dec 13, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  1261da47
11 Dec, 2025 1 commit

wenjh authored Dec 10, 2025


Signed-off-by: wenjh <wenjh@sugon.com>

Mutex group gemm
Signed-off-by: wenjh <wenjh@sugon.com>

do while group gemm
Signed-off-by: wenjh <wenjh@sugon.com>

Remove mutex
Signed-off-by: wenjh <wenjh@sugon.com>

0a90777e

26 Nov, 2025 1 commit
- Use hipFreeAsync rather than hipFree · 772a941a
  wenjh authored Nov 26, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  772a941a
12 Nov, 2025 5 commits
- fix build error · 623ae795
  wenjh authored Nov 12, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  623ae795
- Sync All on groupedgemm. · b3833972
  wenjh authored Nov 12, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  b3833972
- Fix build error · 66bd0b32
  wenjh authored Nov 12, 2025
  
  66bd0b32
- Fix build error · bcaa56b9
  wenjh authored Nov 12, 2025
  
  bcaa56b9
- Fix hipblaslt handle manage · f791516e
  wenjh authored Nov 12, 2025
  
  f791516e
08 Nov, 2025 1 commit
- Fix user args core dump in mt · c1161fb1
  wenjh authored Nov 08, 2025
  
  c1161fb1
16 Oct, 2025 2 commits
- [DCU] remove redundant gemm · 6d43cba8
  yuguo authored Oct 16, 2025
  
  6d43cba8
- [DCU]Fix memory overflow and test-didistributed in L1_pytorch_istributed_unittest · 77e9b318
  tabuchixiangcai3 authored Oct 16, 2025
```
Signed-off-by: Tangao <2205747538@qq.com>
```
  77e9b318
18 Sep, 2025 2 commits
- Adapt to changes of hipblaslt · d81f8119
  wenjh authored Sep 18, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  d81f8119
- [DCU] fix compile · d55fdf7b
  yuguo authored Sep 18, 2025
  
  d55fdf7b
12 Sep, 2025 1 commit
- Fix std::bad_alloc on test_float8blockwisetensor.py · 4922108e
  wenjh authored Sep 12, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  4922108e
11 Sep, 2025 1 commit
- [Perf] blockwise 1d better perf · fca88163
  wenjh authored Sep 11, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  fca88163
02 Sep, 2025 1 commit
- Fix build problems while not support fp4 · 8e0fd518
  wenjh authored Sep 02, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
  8e0fd518
28 Aug, 2025 2 commits
- [DCU] fix quantize bug · d86ee4c8
  yuguo authored Aug 28, 2025
  
  d86ee4c8
- [DCU] fix bugs · 546bb548
  yuguo authored Aug 28, 2025
  
  546bb548
27 Aug, 2025 2 commits

Add cuBLASMp-backed GEMM-like API to TE common (#1824) · d7874aad

Vladimir Cherepanov authored Aug 26, 2025



* Pick up cuBLASMp during build
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Saving...
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Change lib order to fix link error
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Saving...
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Context creation, incomplete...
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Test fixure
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Saving...
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* A sanity AgGemm test, failing...
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Saving...
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Fix axes
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Take care of uneven distribution
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Use MPI to get position of local matrices
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Refactor
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Refactor & fixes
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Saving...
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Gemm-RS
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Gemm-AR, not working...
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Fixes
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Setting all-reduce epilogue for gemm-ar
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Use supported shapes for GEMM-AR
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Tweak tolerance
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* First shot at fp8
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Use TensorHolder in tests
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* More test configs
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Support comm_sm_count
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Parametrize dtypes for A, B and D separately
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Tweak scaling
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Amax ptr
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Flags parity with cublas_gemm, saving...
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Cleanup
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Bias tests
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Fix bias test
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Aux, saving...
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* aux_ld
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* A fix
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Use test::Tensor
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Set scale inv
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Remove unsupported test configs
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Tweak tests
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Replace libcal with NCCL
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Add NVTX markers to API functions
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Tweak GemmAr tests
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* More test config
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

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

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

Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Fix merge fallout
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Remove MPI dependency, comment API, add algo parameter
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Fix nvshmem dependency
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Fix nvshmem build
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Excluse CommGemm tests from L0_cppunittest
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Add cpp_distributed sh file for CI
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Adapt tp TensorAllocator
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

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

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



* Skip GemmAr test on unsupported HW
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Oversibscribe is needed on some clusters
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Fix incomplete libcal removal
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Move CI tests to L1
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Rename context to include NVTE prefix
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Remove leftover code
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* NVTE_WITH_CUBLASMP off by default
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* More detailed NVTE_CHECK diag
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

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

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



* Comment API
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Include stdbool header for legacy C compilers
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Remove now unused argument
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Abstract away cuBLASMp algo behind our own enum
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* More detailed shape diag messages
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

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

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



* Update transformer_engine/common/include/transformer_engine/comm_gemm.h
Co-authored-by: Przemyslaw Tredak <ptrendx@gmail.com>
Signed-off-by: Vladimir Cherepanov <56651474+mk-61@users.noreply.github.com>

* Add license
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

---------
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>
Signed-off-by: Vladimir Cherepanov <56651474+mk-61@users.noreply.github.com>
Co-authored-by: Vladimir Cherepanov <vcherepanov@nvidia.com>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Co-authored-by: Przemyslaw Tredak <ptrendx@gmail.com>

d7874aad

[DCU] fix compile · 5b6190b2
yuguo authored Aug 27, 2025

5b6190b2

26 Aug, 2025 5 commits

Revert "[Common] PDL for Blockwise Quantization" (#2115) · 58c3ac80

jberchtold-nvidia authored Aug 26, 2025

Revert "[Common] PDL for Blockwise Quantization (#2066)"

This reverts commit ebca6153

.
Signed-off-by: Jeremy Berchtold <jberchtold@nvidia.com>

58c3ac80

Bump cuDNN FE to 1.14.0 (#2072) · d2615d1c

vcherepanov-nv authored Aug 26, 2025



* Bump cuDNN FE to 1.14.0
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Change submodule hash
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Pick up a cuDNN FE fix
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* New model configs in tests
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

* Exclude cuDNN backend for some configs
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

---------
Signed-off-by: Vladimir Cherepanov <vcherepanov@nvidia.com>

d2615d1c

Revert "[Common] PDL for Quantization Kernels" (#2114) · 4572dbef

jberchtold-nvidia authored Aug 26, 2025

Revert "[Common] PDL for Quantization Kernels (#2001)"

This reverts commit bfab8c67

.
Signed-off-by: Jeremy Berchtold <jberchtold@nvidia.com>

4572dbef

Fix incorrect version checks for atomic GEMM (#2095) · 3a4136b6

Tim Moon authored Aug 20, 2025



* Fix incorrect version checks for atomic GEMM
Signed-off-by: Tim Moon <tmoon@nvidia.com>

* Fix typo
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>

3a4136b6

[DCU] fix · 4927d10e
yuguo authored Aug 26, 2025

4927d10e

23 Aug, 2025 1 commit
- [DCU] tensorwise int8 gemm surpport bias · 11864d3d
  yuguo authored Aug 23, 2025
  
  11864d3d
21 Aug, 2025 1 commit
- [DCU] tensorwise int8 train opt · 7a923605
  yuguo authored Aug 21, 2025
  
  7a923605
15 Aug, 2025 1 commit

Fuse linear+scale+add (#2042) · c654e4fe

Jan Bielak authored Aug 15, 2025



* Add `nvte_cublas_gemm_scaled`
Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* Support use of `alpha` and `beta` in `tex.generic_gemm`
Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* Support use of `alpha` and `beta` in `general_gemm`
Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* Support use of `alpha` and `beta` in `BasicLinear._functional_forward` and `BasicLinear._functional_backward`
Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* Add `ForwardLinearScaleAdd` fusion
Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* Add `BackwardLinearScale` fusion
Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* Apply suggestions from code review
Signed-off-by: Jan Bielak <jbielak@nvidia.com>

* Remove calls to `validate_gemm_scale` from `BasicLinear`
Signed-off-by: Jan Bielak <jbielak@nvidia.com>

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

c654e4fe

14 Aug, 2025 3 commits

[Core] Add launch bounds to swizzle kernels (#2076) · 12065ac2

Kirthi Shankar Sivamani authored Aug 14, 2025



Add launch bounds to swizzle kernel, use empty scale inv
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>

12065ac2

[Common] Reduce CUDA driver calls (#2067) · c582f6be

Xin Yao authored Aug 14, 2025



* reduce driver calls
Signed-off-by: Xin Yao <xiny@nvidia.com>

* reduce driver calls
Signed-off-by: Xin Yao <xiny@nvidia.com>

* adjust tests to capture this
Signed-off-by: Xin Yao <xiny@nvidia.com>

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

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



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

c582f6be

fix: update grad_output quant to avoid redundant work (#1736) · 44fbe9e6

Kshiteej K authored Aug 14, 2025



* fix: update grad_output quant to avoid redundant work
Signed-off-by: kshitij12345 <kshitijkalambarkar@gmail.com>

* add test
Signed-off-by: kshitij12345 <kshitijkalambarkar@gmail.com>

* don't keep only columnwise quant if requires_dgrad=False
Signed-off-by: kshitij12345 <kshitijkalambarkar@gmail.com>

* fix stray merge
Signed-off-by: kshitij12345 <kshitijkalambarkar@gmail.com>

* fix for ctx.use_bias is True case
Signed-off-by: kshitij12345 <kshitijkalambarkar@gmail.com>

* Skip if FP8 not available
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>

---------
Signed-off-by: kshitij12345 <kshitijkalambarkar@gmail.com>
Signed-off-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>
Co-authored-by: Kirthi Shankar Sivamani <ksivamani@nvidia.com>

44fbe9e6

13 Aug, 2025 3 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

[Common] PDL for Blockwise Quantization (#2066) · ebca6153

Xin Yao authored Aug 13, 2025



* enable PDL for blockwise qunatization kernels
Signed-off-by: Xin Yao <xiny@nvidia.com>

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

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

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



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

ebca6153

12 Aug, 2025 2 commits

[PyTorch] Fix normalization+amax forward CS fusion to work for untuned kernels (#2061) · 05d3b7b5

Jan Bielak authored Aug 12, 2025



* Compute amax in normalization forward in current scaling in untuned kernels
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: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>

05d3b7b5

Use hipMemcpy rather than hipMemcpyAsync · 1f97aebb
wenjh authored Aug 12, 2025
```
Signed-off-by: wenjh <wenjh@sugon.com>
```
1f97aebb