v2.4.17

What's Changed

• [NMS] Add nms f32 cuda kernel. by @bear-zd in https://github.com/DefTruth/CUDA-Learn-Notes/pull/102
• [HGEMM] Add some note to collective store by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/103
• [HGEMM] Add HGEMM MMA Col Major Kernel by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/104
• [HGEMM] Update HGEMM benchmark scripts by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/105
• [HGEMM] Add Warp Swizzle as template param by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/106
• [HGEMM] add -Xptxas -v compile flag by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/107
• [HGEMM] Try reduce registers usage by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/108
• [HGEMM] Update HGEMM MMA/WMMA Usage by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/109
• [HGEMM][Docs] Add HGEMM Supported Matrix by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/110
• [HGEMM] Add M=N=K option for benchmark by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/111
• [HGEMM] Update HGEMM/SGEMM Supported Matrix by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/112
• [README] Update HGEMM/SGEMM Supported matrix by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/113
• [Docs] Update HGEMM/SGEMM Supported Matrix by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/114

Full Changelog: DefTruth/CUDA-Learn-Notes@v2.4.16...v2.4.17

HGEMM Warp Swizzle/Reg Buffers

What's Changed

• [HGEMM] HGEMM MMA with Reg Double Buffers by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/99
• [HGEMM] ldmatrix.x4.trans with reg double buffers by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/100
• [HGEMM] collective store via warp shfl&reg reuse by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/101

Full Changelog: DefTruth/CUDA-Learn-Notes@v2.4.15...v2.4.16

HGEMM Up to 115 TFLOPS:L20

What's Changed

• [HGEMM] Add MMA 16816 swizzle, Up to 115 TFLOPS by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/98

Full Changelog: DefTruth/CUDA-Learn-Notes@v2.4.13...v2.4.15

HGEMM Up to 113 TFLOPS:L20

What's Changed

• [Mat][Trans] Add f32/f32x4 row/col first kernel by @bear-zd in https://github.com/DefTruth/CUDA-Learn-Notes/pull/89
• [Docs][Contribute] Add How to contribute Notes by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/90
• [HGEMM] optimize SMEM padding, up to 113 TFLOPS by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/92
• [Mat][Trans] Add f32x4_shared/bcf row/col first kernel. by @bear-zd in https://github.com/DefTruth/CUDA-Learn-Notes/pull/91
• [Docs] rename mat_transpose -> mat-transpose by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/93
• [HGEMM] Add GeForce RTX 3080 Laptop benchmark by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/94
• [HGEMM] update HGEMM benchmark option by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/95
• [HGEMM] Refactor HGEMM WMMA 161616 kernels by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/96
• [HGEMM] Update HGEMM WMMA Benchmark by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/97

Full Changelog: DefTruth/CUDA-Learn-Notes@v2.4.12...v2.4.13

v2.4.12 SGEMM TF32 Swizzle

What's Changed

• [SGEMM] SGEMM TF32 Thread Block Swizzle by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/84
• [HGEMM] mma4x4_warp4x4_stages with swizzle by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/86
• [SWISH] support Swish F32/F16 kernel by @wangzijian1010 in https://github.com/DefTruth/CUDA-Learn-Notes/pull/85
• [SGEMM] Update SGEMM TF32 Benchmark by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/87

New Contributors

• @wangzijian1010 made their first contribution in https://github.com/DefTruth/CUDA-Learn-Notes/pull/85

Full Changelog: DefTruth/CUDA-Learn-Notes@v2.4.11...v2.4.12

v2.4.11 HGEMM Block Swizzle

What's Changed

• [Docs] Update README.md by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/81
• [HEGMM] HGEMM WMMA Thread Block Swizzle by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/82
• [HGEMM] make thread block swizzle stride as N/4 by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/83

Full Changelog: DefTruth/CUDA-Learn-Notes@v2.4.10...v2.4.11

v2.4.10 SGEMM TF32 Stage 2/3

What's Changed

• [HGEMM] HGEMM WMMA Stage mma4x2+warp4x4 by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/76
• [SGEMM] Add SGEMM WMMA TF32 Stage2/3 by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/77
• [SGEMM] Add cuBLAS SGEMM F32/TF32 baseline by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/78
• [SGEMM] Add Kernel cudaFuncSetAttribute hint by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/79
• [RoPE] Add minimal RoPE f32/f32x4 pack impl by @bear-zd in https://github.com/DefTruth/CUDA-Learn-Notes/pull/80

Full Changelog: DefTruth/CUDA-Learn-Notes@v2.4.9...v2.4.10

v2.4.9 HGEMM WMMA Stage

What's Changed

• [HGEMM] Add HGEMM WMMA Double Buffers by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/69
• [Embedding] Add embedding kernel f32/x4/x4_pack, f16/x8/x8_pack by @bear-zd in https://github.com/DefTruth/CUDA-Learn-Notes/pull/68
• [HGEMM] Add HGEMM mma4x2, warp2x4x2 kernel by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/70
• [HGEMM] HGEMM WMMA with Reg double buffers by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/71
• [HGEMM] Add HGEMM WMMA Stage 3/4 Kernel by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/74
• [Softmax] Add online softmax f32x4 pack kernel by @bear-zd in https://github.com/DefTruth/CUDA-Learn-Notes/pull/73
• [HEGMM][Bugfix] fix HGEMM Stage cp.async error by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/75

Full Changelog: DefTruth/CUDA-Learn-Notes@v2.4.8...v2.4.9

v2.4.8 HGEMM WMMA Part-1

What's Changed

• [GELU] Add f32/x4, f16/x2/x8/x8pack kernel. by @bear-zd in https://github.com/DefTruth/CUDA-Learn-Notes/pull/66
• [HGEMM] HGEMM Tensor Cores Support Part-1 by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/67

Full Changelog: DefTruth/CUDA-Learn-Notes@v2.4.7...v2.4.8

v2.4.7 SGEMM Copy Async

What's Changed

• [SGEMM][Async] Add naive copy async SGEMM by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/64
• [SGEMM][Async] Add K16 + Copy Async Kernel by @DefTruth in https://github.com/DefTruth/CUDA-Learn-Notes/pull/65

Full Changelog: DefTruth/CUDA-Learn-Notes@v2.4.6...v2.4.7