NVIDIA/cutlass
1
0
Fork 0
mirror of https://github.com/NVIDIA/cutlass.git synced 2026-05-12 09:15:56 +00:00
Code Issues Packages Projects Releases Wiki Activity
831 Commits 21 Branches 52 Tags
2aedca6f5eea09a6928cbfc78de926a37d5f765d
Commit Graph

831 Commits

This Branch
This Branch
All Branches
Author SHA1 Message Date
Haicheng Wu
2aedca6f5e Bump CUTLASS version to 4.4.0 2026-02-25 00:01:56 -05:00
Haicheng Wu
6450964b57 Update README 2026-02-24 23:55:55 -05:00
Haicheng Wu
284449fa5b Revise chagnelog 2026-02-24 23:54:56 -05:00
Haicheng Wu
0853d81d70 Revise README 2026-02-24 15:32:17 -05:00
Linfeng Zheng
3476ddb7bd remove mixed_input_fmha_prefill (#3041) 2026-02-18 07:59:01 -05:00
Yihan Chen
291300ffff [CuTeDSL] implment a cta-level norm example (both layernorm and rmsnorm) (#3009) 
* kernel impl

* add copyright
 2026-02-14 17:54:03 +08:00
aragorn-guan
f9a5f76b7a Replace fence proxy to the latest routine code in examples/distributed/all_reduce_tma.py (#3027) 2026-02-14 17:51:20 +08:00
drazi
ec7e6cb17b Merge pull request #2971 from rsmallblue/tvm-ffi 
[CuTeDSL]fix tvm-ffi path in from_dlpack
 2026-02-14 14:14:10 +08:00
Yuan Xiaolan
395ab575f6 Merge branch 'main' into tvm-ffi 2026-02-14 13:35:28 +08:00
Junkai-Wu
d4bbf728ca v4.4 tag release update. (#3032) 2026-02-13 23:27:58 -05:00
drazi
01687cfba1 Merge pull request #3004 from tridao/add-sub-packed-f32x2 
[CuTeDSL] Add sub_packed_f32x2 operation
 2026-02-13 20:46:26 +08:00
drazi
5c42d0f28c Merge pull request #3021 from tridao/clc_no_multicast 
[Cute-DSL] Add option for issue_clc_query without multicast
 2026-02-13 20:45:52 +08:00
drazi
1d36152f34 Merge pull request #3022 from tridao/nvvm_fmin 
[Cute-DSL] Add cute.arch.fmin by calling nvvm
 2026-02-13 20:45:08 +08:00
Tri Dao
244e8d00d5 [Cute-DSL] Add cute.arch.fmin by calling nvvm 2026-02-11 14:23:09 -05:00
Tri Dao
5b83b34afd [Cute-DSL] Add option for issue_clc_query without multicast 2026-02-11 14:19:29 -05:00
aragorn-guan
8dbce01473 [CuTeDSL] Distributed example, using TMA load to access remote memory rank-by-rank, reducing in cta, broadcast result to all ranks by multimem TMA store (#2970) 2026-02-11 11:54:00 +08:00
drazi
71aa7a0abc Merge pull request #2919 from pbelevich/patch-1 
Refactor binary_op functions to remove unused result parameter
 2026-02-11 11:48:58 +08:00
Tri Dao
51935551fb [CuTeDSL] Add sub_packed_f32x2 operation 
Add subtraction operation for packed f32x2 values, following the same
pattern as the existing add_packed_f32x2 and mul_packed_f32x2 operations.

Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
 2026-02-04 21:18:46 +07:00
Junkai-Wu
6b3e607b85 v4.4 release update v2. (#2999) 2026-02-03 20:48:31 -05:00
yuanxiaolan
de161925a5 pass in stream=-1 2026-02-03 11:59:14 +08:00
yuanxiaolan
de198b2419 fix tvm-ffi path in from_dlpack 2026-02-03 11:59:13 +08:00
Hua Huang
1cfbb53a23 [CuTeDSL] Fix: SM100 block-scale gemm overlapping accumulator (#2995) 
* Fix: SM100 block-scale gemm overlapping accumulator

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

* Also include threads_per_warp fix

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

---------

Signed-off-by: Hua Huang <huah@nvidia.com>
 2026-02-03 11:01:41 +08:00
dongxiao
a4eb0e05f6 fix performance inssues in cute-dsl examples for 4.4-ctk13.1 release (#2988) 
* fix grouped gemm

* fix mixed input gemm

* fix mixed input grouped gemm

* fix version checking

* use advanced compiler options

* fix comment

* rename advanced compiler configs to adcanced compiler control

* fix comment

* fix name

* fix name
 2026-01-30 13:31:04 +08:00
myu-guo
d252b01300 fix performance regression in cute-dsl examples for 4.4-ctk13.1 release (#2990) 
* fix regression with cu13.1

* update
 2026-01-30 13:30:49 +08:00
Xiao Song
acb45938e9 Update nvvm API call from nvvm enum to str (#2985) 2026-01-27 17:28:29 +08:00
Xiao Song
7a14467776 update api usage (#2969) 2026-01-27 15:33:22 +08:00
drazi
51f82812ec Merge pull request #2891 from ColinPeppler/main 
docs: note when DSL dumps are populated
 2026-01-26 17:38:27 -08:00
Junkai-Wu
9fba3195f9 v4.4 update. (#2979) 2026-01-24 11:46:17 -05:00
Qi Yuhang
2fafefb7b9 [Bug Fix]Set NumSplitsM to 1 when TileShapeM < 128 in sm90 fp8 blockwise scaling CollectiveMma (#2965) 
* Fix NumSplitsM when TileShapeM < 128.

* Use cute::conditional_t to replace std::conditional_t.
 2026-01-23 15:56:52 +08:00
Johnsonms
0edaa6e47d Fix out-of-bounds TMA access in wgmma_tma_sm90 tutorial (#2945) 2026-01-23 12:54:12 +08:00
Aidan Do
431d070fcb [docs] Add additional tip for generating less kernels in blockwise (#2940) 
- Running without this generates a lot of kernels
- Clarified CMake configuration for selecting GEMM kernels and added details on kernel generation granularity.
 2026-01-23 12:53:51 +08:00
Qi Yuhang
667446a9dd [Doc]Fix Mode Name and Stride in 0t_mma_atom.md (#2910) 
* Fix wrong stride in 0t_mma_atom.md

* Fix mode name.

* Update media/docs/cpp/cute/0t_mma_atom.md

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

---------

Co-authored-by: Cris Cecka <ccecka@users.noreply.github.com>
 2026-01-23 12:53:30 +08:00
Aidan Do
3f5bafb326 [Cutlass profiler] Fix SM100 FP8 nosmem epilogue shape_div “Divisibility Condition” for non‑multiple‑of‑64 N tiles (#2946) 
* .

* .

* .

* .

* .

* .

* .
 2026-01-20 15:27:34 +08:00
Tianqi Chen
1e6da09275 [DOCS] Update docs to precisely describe env stream scenario (#2824) 
Since the term synchronize may cause confusion that user thought it means stream sync, it simply means we pass the right current stream as env stream
 2026-01-20 09:16:37 +08:00
Benjamin Leff
8debf77437 fix: 2305 omissions (#2957) 2026-01-14 00:55:05 -05:00
Brian K. Ryu
147f5673d0 New RMS Norm example with unit tests (#2917) 
* Add rmsnorm example

* Address reviewer comments. (1) use the cute.runtime definition directly. (2) use the nvvm_wrapper's warp reduce directly

* Separate out reduce.py

* Change copyright notice years
 2026-01-13 09:05:31 +08:00
Johnsonms
8c52459504 Fix incorrect tensor layout strides in Blackwell MMA tutorial comments (#2921) 2026-01-09 01:02:41 -05:00
kf-zhang
0deda34b9f fix typo (#2884) 2026-01-09 00:57:06 -05:00
Junkai-Wu
0d2b201e8c v4.3.5 update. (#2934) 
* v4.3.5 update.

* Update copyright to 2026
 2026-01-08 15:02:56 -05:00
Wenxuan Tan
f86feb0aa8 Fix idx2crd docstring (#2914) 
* fix idx2crd docstring

* fix

* fix
 2026-01-07 13:11:38 -05:00
Andrew Yooeun Chun
eb61c91147 Fix CUDA version checking in examples (#2894) 
* examples: update CUDA version requirements in Blackwell examples

* examples: fix comments to specify the correct CUDA version requirement
 2026-01-07 00:20:37 -05:00
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
Pavel Belevich
b6d7703e02 Refactor binary_op functions to remove unused result parameter 2026-01-02 11:23:43 -05:00
Pavel Belevich
f9bedd9096 Fix print statement for floor division result 2026-01-02 11:15:15 -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