Commit Graph

787 Commits

Author SHA1 Message Date
Aidan Do
670480df3a Fix SFB Layout scale granularity representation (#2924) 2026-01-06 23:55:21 -05:00
veritas-Qiu
61b560983a remove useless line (#2926)
the parameter workspace is marked as unused like other kernels, but it is actually used after 3.3.0, so the code which mark it as unused could be removed.
2026-01-06 23:54:08 -05:00
dePaul Miller
7127592069 Replace CUDA driver API with runtime API (#2928)
Co-authored-by: dePaul Miller <23461061+depaulmillz@users.noreply.github.com>
2026-01-05 13:50:44 -05:00
questa-quan-wang
2aee73922c Minor fix for testing of blockscaled dense GEMM with TMA prefetch (#2930)
* new example with TMA prefetch feature targeting for DRAM latency bound cases

* minor fix to resitrct as 100a arch

* typo

* apply arch for whole pytest

---------

Co-authored-by: Questa Wang <questaw@computelab-frontend-7.nvidia.com>
Co-authored-by: Questa Wang <questaw@umbriel-b200-145.ipp4a1.colossus.nvidia.com>
2026-01-05 16:36:03 +08:00
tsu-bin
3d9de19bb7 add constexpr specifier to make_tiled_copy (#2875) 2026-01-03 15:39:43 -05:00
Haicheng Wu
853ad93d60 Update README.md 2025-12-24 00:21:59 -05:00
Haicheng Wu
34a81f0497 Update driver bug workaround description in CHANGELOG
Clarified the description of a driver bug workaround in the CHANGELOG.
2025-12-24 00:20:21 -05:00
questa-quan-wang
3f4c086d09 new example with TMA prefetch feature targeting for DRAM latency bound cases (#2881)
Co-authored-by: Questa Wang <questaw@computelab-frontend-7.nvidia.com>
2025-12-23 15:29:48 +08:00
Junkai-Wu
b7ecaa605d v4.3.4 update v2. (#2898) 2025-12-22 22:28:26 -05:00
Junkai-Wu
7f5fe3edf1 v4.3.4 update. (#2892) 2025-12-21 11:49:12 -05:00
dongxiao
331e2f451c add missing condition for sync (#2889) 2025-12-19 11:00:30 +08:00
Qi Yuhang
ebf3165efb [Bug Fix]Bypass launch grids for SM120 Kernel with SM90 Mainloop & SM100 TileScheduler (#2865)
* Delete unused #ifdef/#endif. Bypass sm120 case.

* Add todo.

* Fix pingpong.

* Revert "Add todo."

This reverts commit 246cb42091b1ed1b89c1eac2312eed6625a754fd.

* Refine name.

Refine name again.

* Apply suggestions from code review

Skip `is_last_tile` for all sm120 kernels.

Co-authored-by: Junkai-Wu <junkaiw@nvidia.com>

* Skip early stop for sm120 kernel.

* Fix typo.

---------

Co-authored-by: Junkai-Wu <junkaiw@nvidia.com>
2025-12-18 08:51:38 +08:00
Haicheng Wu
d4e16f5d4e Bump version from 4.2.1 to 4.3.3 2025-12-11 23:58:38 -05:00
Junkai-Wu
d3a5492381 v4.3.3 update. (#2868) 2025-12-11 00:26:58 -05:00
Amin Sedaghat
49bd6bf1ba fix print_layout printf format in device code (#2688)
* fix print_layout printf format in device code

* Replace %.*s format specifier with explicit loop
* Remove unused delim variable

The printf format %.*s with dynamic width does not work correctly
in CUDA device code, causing literal %.*s to appear in output.

Fixes #2496

* Update include/cute/util/print_tensor.hpp

Co-authored-by: Cris Cecka <ccecka@users.noreply.github.com>

* Update include/cute/util/print_tensor.hpp

Co-authored-by: Cris Cecka <ccecka@users.noreply.github.com>

---------

Co-authored-by: Cris Cecka <ccecka@users.noreply.github.com>
2025-12-10 08:57:56 +08:00
Junkai-Wu
c6dfdf8375 Merge pull request #2719 from HydraQYH/dev_support_pdl_for_sm90_gemm_array_tma_ws
Support PDL for SM90 Array TMA GEMM
2025-12-09 18:32:15 +08:00
HydraQYH
95f8beb44c Revert "Remove unnecessary #ifdef #endif for general gemm."
This reverts commit 17ffd56dfe.
2025-12-09 11:52:23 +08:00
HydraQYH
17ffd56dfe Remove unnecessary #ifdef #endif for general gemm. 2025-12-06 10:20:35 +08:00
HydraQYH
ff7f2dcdfb Remove duplicated cutlass::arch::wait_on_dependent_grids(); 2025-12-06 10:20:35 +08:00
HydraQYH
929e1e0259 Remove unnecessary #ifdef / #endif for launch_dependent_grids. 2025-12-06 10:20:35 +08:00
HydraQYH
b6ad6db219 Delete unnecessary #ifdef / #endif. 2025-12-06 10:20:35 +08:00
HydraQYH
e1b2ec57e3 Hoist waits above the warp specialized region. 2025-12-06 10:20:35 +08:00
HydraQYH
1e5f95cbbe Support PDL in sm90_gemm_array_tma_warpspecialized_cooperative 2025-12-06 10:20:35 +08:00
HydraQYH
acf5990cc2 Refine position for wait_on_dependent_grids. 2025-12-06 10:20:35 +08:00
HydraQYH
91de7891a5 Support PDL in sm90_gemm_array_tma_warpspecialized_pingpong.hpp 2025-12-06 10:20:35 +08:00
Haicheng Wu
1e7d3e030b Update CHANGELOG for CuTe DSL enhancements
Added new environment variable for CuTe DSL cache directory.
2025-12-05 13:49:57 -05:00
Haicheng Wu
c4744f706e Bump version from 4.2.1 to 4.3.2 2025-12-05 13:45:16 -05:00
Linfeng Zheng
f6402fcd5e add pytest support for tutorial gemm (#2826)
* add pytest support for tutorial gemm

* add license
2025-12-05 08:45:01 -05:00
bangyu shen
7252a2d17e remove internal comment (#2841)
Co-authored-by: bangyus <bangyus@nvidia.com>
2025-12-04 10:36:21 -05:00
Junkai-Wu
bc680c7f67 v4.3.2 update. (#2839) 2025-12-04 10:14:32 -05:00
bangyu shen
52ae719eda [examples][CuTeDSL] init commit for distirbuted examples (#2806)
* init commit for distirbuted examples

* better OOB protection

* and try import to nvshmem for better error message and a READMME.md to introduce nvshmem and multimem instructions

* add some lamport explanation

* enhance f8 output and warn that f8 output can have nan in it

* tell user why we need complicate data conversions in ref check part

* tell user we don't support nvshmem device function

---------

Co-authored-by: bangyus <bangyus@nvidia.com>
2025-12-01 22:25:40 -05:00
Haicheng Wu
5e847d37c4 Bump version from 4.2.1 to 4.3.1 2025-12-01 22:13:19 -05:00
Haicheng Wu
f16068b4db Bump version from 4.2.0 to 4.3.1 2025-12-01 22:12:20 -05:00
Haicheng Wu
1acfe141af Bump version from 4.2.1 to 4.3.1 2025-12-01 22:11:13 -05:00
Haicheng Wu
f11375bf91 Bump CUTLASS patch version to 1 2025-12-01 22:08:52 -05:00
Shreya Gaur
af8d5dfa54 bug fix for example 92 (#2830)
Co-authored-by: Shreya Gaur <shgaur@dc2-container-xterm-012.prd.it.nvidia.com>
Co-authored-by: Shreya Gaur <shgaur@2u2g-spr-0015.ipp4a1.colossus.nvidia.com>
2025-12-01 22:02:59 -05:00
drazi
ec8daf642d Merge pull request #2809 from whatdhack/patch-1
Update notebook title from 'Tour to' to 'Tour of'
2025-11-28 18:07:34 +08:00
drazi
5016493cc0 Merge pull request #2813 from fengxie/ftse/fix/example
Refactor TVM FFI examples and update doc
2025-11-28 09:07:15 +08:00
Fung Xie
8588d099e4 refactored doc 2025-11-27 17:04:20 -08:00
Fung Xie
8fc9bc5dda update doc 2025-11-27 17:03:51 -08:00
Fung Xie
f71892b824 update doc 2025-11-27 17:03:03 -08:00
Fung Xie
03aa211310 update doc 2025-11-27 17:02:59 -08:00
Fung Xie
286781a1fb add requirements.txt 2025-11-27 17:02:27 -08:00
Fung Xie
2664cac685 enhanced the example for tvm-ffi 2025-11-27 17:02:26 -08:00
Fung Xie
b9154d65b3 update examples for tvm-ffi 2025-11-27 17:02:26 -08:00
Fung Xie
afe2f71522 reorganize examples for tvm-ffi 2025-11-27 17:02:26 -08:00
Fung Xie
739fffce27 fix TVM FFI doc and update example 2025-11-27 17:02:26 -08:00
Junkai-Wu
1de3a576cc v4.3.1 update. (#2817) 2025-11-27 09:49:30 -05:00
Shreya Gaur
2052fd3885 Blockscaled Ragged Contiguous Grouped Gemm for MoEs (#2790)
* Adding blockscaled ragged contiguous grouped gemm for MoEs

* cleaning up the example

* introduction to example improved

---------

Co-authored-by: Shreya Gaur <shgaur@dc2-container-xterm-012.prd.it.nvidia.com>
2025-11-26 20:16:49 -05:00
whatdhack
4a55379686 Update notebook title from 'Tour to' to 'Tour of'
Grammar check . LLM's can quickly clean up such issues.
2025-11-24 20:11:14 -08:00