From c10770b8c7d6b391720e3d3bcd2874ba874aab83 Mon Sep 17 00:00:00 2001
From: Hui Zhou <90194592+kitecats@users.noreply.github.com>
Date: Mon, 2 Jun 2025 17:35:56 +0800
Subject: [PATCH 01/57] Use 128-bit data loading (#326)

---
 kernels/flash-attn/cutlass/flash_attn_cute.cu | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/kernels/flash-attn/cutlass/flash_attn_cute.cu b/kernels/flash-attn/cutlass/flash_attn_cute.cu
index c570910f..6fe03fa0 100644
--- a/kernels/flash-attn/cutlass/flash_attn_cute.cu
+++ b/kernels/flash-attn/cutlass/flash_attn_cute.cu
@@ -37,7 +37,7 @@ struct FlashAttnConfig {
 
   // Gmem2Smem config
   using GmemCopyAtom =
-      Copy_Atom<AutoVectorizingCopyWithAssumedAlignment<sizeof(uint128_t)>, T>;
+      Copy_Atom<AutoVectorizingCopyWithAssumedAlignment<sizeof(uint128_t)*8>, T>;
   static constexpr int GmemValsPerLoad = sizeof(uint128_t) / sizeof(T);
   static constexpr int GmemThreadsPerRow =
       HeadDim / GmemValsPerLoad; // each thread reads 128 bit
@@ -56,7 +56,7 @@ struct FlashAttnConfig {
   using SmemCopyAtomTransposed =
       Copy_Atom<SM75_U16x8_LDSM_T, T>; // for column major load
   using SmemCopyAtomO =
-      Copy_Atom<AutoVectorizingCopyWithAssumedAlignment<sizeof(uint128_t)>,
+      Copy_Atom<AutoVectorizingCopyWithAssumedAlignment<sizeof(uint128_t)*8>,
                 T>; // NOTE: stmatrix is only available after sm90, we use a
                     // vectorized copy instead
 

From 0fc44ad465ab9a4a2a2d1cc97706d5d6e90ac7b7 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Mon, 2 Jun 2025 21:23:35 +0800
Subject: [PATCH 02/57] Create FUNDING.yml (#327)

---
 .github/FUNDING.yml | 15 +++++++++++++++
 1 file changed, 15 insertions(+)
 create mode 100644 .github/FUNDING.yml

diff --git a/.github/FUNDING.yml b/.github/FUNDING.yml
new file mode 100644
index 00000000..0bab9ea8
--- /dev/null
+++ b/.github/FUNDING.yml
@@ -0,0 +1,15 @@
+# These are supported funding model platforms
+
+github: # Replace with up to 4 GitHub Sponsors-enabled usernames e.g., [user1, user2]
+patreon: # Replace with a single Patreon username
+open_collective: https://opencollective.com/leetcuda # Replace with a single Open Collective username
+ko_fi: # Replace with a single Ko-fi username
+tidelift: # Replace with a single Tidelift platform-name/package-name e.g., npm/babel
+community_bridge: # Replace with a single Community Bridge project-name e.g., cloud-foundry
+liberapay: # Replace with a single Liberapay username
+issuehunt: # Replace with a single IssueHunt username
+lfx_crowdfunding: # Replace with a single LFX Crowdfunding project-name e.g., cloud-foundry
+polar: # Replace with a single Polar username
+buy_me_a_coffee: # Replace with a single Buy Me a Coffee username
+thanks_dev: # Replace with a single thanks.dev username
+custom: # Replace with up to 4 custom sponsorship URLs e.g., ['link1', 'link2']

From b25686855958ab54213469d26a069dc248e7259c Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Mon, 2 Jun 2025 21:24:37 +0800
Subject: [PATCH 03/57] Update FUNDING.yml

---
 .github/FUNDING.yml | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/.github/FUNDING.yml b/.github/FUNDING.yml
index 0bab9ea8..2be6087c 100644
--- a/.github/FUNDING.yml
+++ b/.github/FUNDING.yml
@@ -2,7 +2,7 @@
 
 github: # Replace with up to 4 GitHub Sponsors-enabled usernames e.g., [user1, user2]
 patreon: # Replace with a single Patreon username
-open_collective: https://opencollective.com/leetcuda # Replace with a single Open Collective username
+open_collective: leetcuda # Replace with a single Open Collective username
 ko_fi: # Replace with a single Ko-fi username
 tidelift: # Replace with a single Tidelift platform-name/package-name e.g., npm/babel
 community_bridge: # Replace with a single Community Bridge project-name e.g., cloud-foundry

From 0b5cc1d433814783b624f59ff840d5e9701620fb Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Mon, 2 Jun 2025 21:30:06 +0800
Subject: [PATCH 04/57] Add open-collective badge (#328)

---
 README.md | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/README.md b/README.md
index 245759d5..515eb743 100644
--- a/README.md
+++ b/README.md
@@ -30,9 +30,11 @@
 <div align="center">
   <p align="center">
     <a href="#contribute">🔥🔥 PR Welcome: Add Your Kernel to LeetCUDA! Let's make it Awesome together! 🎉🎉</a>
+    <img src=https://opencollective.com/leetcuda/contributors.svg>
   </p>
 </div>
 
+
 ## 📖 News 🔥🔥
 <div id="news"></div>
 

From ea97bb497c63d369177befe7c52bbbd7d77adfb2 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Mon, 2 Jun 2025 21:36:01 +0800
Subject: [PATCH 05/57] Update open-collective contributors badge (#329)

---
 README.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/README.md b/README.md
index 515eb743..6a31742b 100644
--- a/README.md
+++ b/README.md
@@ -30,7 +30,7 @@
 <div align="center">
   <p align="center">
     <a href="#contribute">🔥🔥 PR Welcome: Add Your Kernel to LeetCUDA! Let's make it Awesome together! 🎉🎉</a>
-    <img src=https://opencollective.com/leetcuda/contributors.svg>
+    <a href=https://github.com/xlite-dev/LeetCUDA/graphs/contributors > <img src=https://opencollective.com/leetcuda/contributors.svg height=40px > </a>
   </p>
 </div>
 

From edd2a5b910c563023a9b9a1b22edb298f959a902 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Mon, 2 Jun 2025 21:51:10 +0800
Subject: [PATCH 06/57] Update FUNDING.yml

---
 .github/FUNDING.yml | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/.github/FUNDING.yml b/.github/FUNDING.yml
index 2be6087c..b4c810a6 100644
--- a/.github/FUNDING.yml
+++ b/.github/FUNDING.yml
@@ -1,6 +1,6 @@
 # These are supported funding model platforms
 
-github: # Replace with up to 4 GitHub Sponsors-enabled usernames e.g., [user1, user2]
+github: xlite-dev # Replace with up to 4 GitHub Sponsors-enabled usernames e.g., [user1, user2]
 patreon: # Replace with a single Patreon username
 open_collective: leetcuda # Replace with a single Open Collective username
 ko_fi: # Replace with a single Ko-fi username

From 04cc34db4b77d5e050d505a8fca840abb9fc8f51 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Mon, 2 Jun 2025 21:52:44 +0800
Subject: [PATCH 07/57] Update FUNDING.yml

---
 .github/FUNDING.yml | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/.github/FUNDING.yml b/.github/FUNDING.yml
index b4c810a6..2be6087c 100644
--- a/.github/FUNDING.yml
+++ b/.github/FUNDING.yml
@@ -1,6 +1,6 @@
 # These are supported funding model platforms
 
-github: xlite-dev # Replace with up to 4 GitHub Sponsors-enabled usernames e.g., [user1, user2]
+github: # Replace with up to 4 GitHub Sponsors-enabled usernames e.g., [user1, user2]
 patreon: # Replace with a single Patreon username
 open_collective: leetcuda # Replace with a single Open Collective username
 ko_fi: # Replace with a single Ko-fi username

From 45d1112e829f2abc7772729434466f6d446da108 Mon Sep 17 00:00:00 2001
From: Hui Zhou <90194592+kitecats@users.noreply.github.com>
Date: Fri, 6 Jun 2025 21:02:08 +0800
Subject: [PATCH 08/57] feat: add cute hgemv implement (#331)

* Use 128-bit data loading

* add a cute hgemv implement

* updata readme

* updata readme

* updata readme
---
 kernels/flash-attn/cutlass/flash_attn_cute.cu |   5 +-
 kernels/hgemv/README.md                       |  20 +-
 kernels/hgemv/hgemv.cu                        |   7 +
 kernels/hgemv/hgemv.py                        |  14 +-
 kernels/hgemv/hgemv_cute.cu                   | 408 ++++++++++++++++++
 5 files changed, 446 insertions(+), 8 deletions(-)
 create mode 100644 kernels/hgemv/hgemv_cute.cu

diff --git a/kernels/flash-attn/cutlass/flash_attn_cute.cu b/kernels/flash-attn/cutlass/flash_attn_cute.cu
index 6fe03fa0..ad2d93b6 100644
--- a/kernels/flash-attn/cutlass/flash_attn_cute.cu
+++ b/kernels/flash-attn/cutlass/flash_attn_cute.cu
@@ -37,7 +37,8 @@ struct FlashAttnConfig {
 
   // Gmem2Smem config
   using GmemCopyAtom =
-      Copy_Atom<AutoVectorizingCopyWithAssumedAlignment<sizeof(uint128_t)*8>, T>;
+      Copy_Atom<AutoVectorizingCopyWithAssumedAlignment<sizeof(uint128_t) * 8>,
+                T>;
   static constexpr int GmemValsPerLoad = sizeof(uint128_t) / sizeof(T);
   static constexpr int GmemThreadsPerRow =
       HeadDim / GmemValsPerLoad; // each thread reads 128 bit
@@ -56,7 +57,7 @@ struct FlashAttnConfig {
   using SmemCopyAtomTransposed =
       Copy_Atom<SM75_U16x8_LDSM_T, T>; // for column major load
   using SmemCopyAtomO =
-      Copy_Atom<AutoVectorizingCopyWithAssumedAlignment<sizeof(uint128_t)*8>,
+      Copy_Atom<AutoVectorizingCopyWithAssumedAlignment<sizeof(uint128_t) * 8>,
                 T>; // NOTE: stmatrix is only available after sm90, we use a
                     // vectorized copy instead
 
diff --git a/kernels/hgemv/README.md b/kernels/hgemv/README.md
index 35526963..a3bda70e 100755
--- a/kernels/hgemv/README.md
+++ b/kernels/hgemv/README.md
@@ -7,8 +7,12 @@
 - [X] hgemv_k32_f16_kernel
 - [X] hgemv_k128_f16x4_kernel
 - [X] hgemv_k16_f16_kernel
+- [X] hgemv_f16_cute_kernel
+- [X] hgemv_f16x8_cute_kernel
+- [X] hgemv_tensor_core_cute_kernel
 - [X] PyTorch bindings
 
+
 ## 测试
 
 ```bash
@@ -21,11 +25,17 @@ python3 hgemv.py
 
 ```bash
 --------------------------------------------------------------------------------
-   out_k32f16: [-0.00585938, -1.91015625, 10.8281250], time:0.00373125ms
-out_k128f16x4: [-0.01171875, -1.89648438, 10.8203125], time:0.00373244ms
-   out_f16_th: [-0.00727463, -1.89648438, 10.8281250], time:0.00845909ms
+   out_k32f16: [15.609375, 2.15234375, -10.9296875], time:0.00324011ms
+out_k128f16x4: [15.609375, 2.15625, -10.9296875], time:0.00322700ms
+out_hgemv_f16_cute: [15.609375, 2.15234375, -10.9296875], time:0.00318646ms
+out_hgemv_f16x8_cute: [15.609375, 2.16015625, -10.9375], time:0.00323176ms
+out_hgemv_tensor_core_cute: [15.6171875, 2.15625, -10.9375], time:0.00531912ms
+   out_f16_th: [15.6171875, 2.15429688, -10.9375], time:0.00889659ms
 --------------------------------------------------------------------------------
-   out_k16f16: [6.03125, -2.15234375, 3.140625], time:0.00377536ms
-   out_f16_th: [6.03125, -2.15234375, 3.140625], time:0.00835538ms
+   out_k16f16: [-6.69140625, -7.2265625, -6.4921875], time:0.00339985ms
+out_hgemv_f16_cute: [-6.69140625, -7.2265625, -6.4921875], time:0.00323296ms
+out_hgemv_f16x8_cute: [-6.6875, -7.2265625, -6.4921875], time:0.00319839ms
+out_hgemv_tensor_core_cute: [-6.6875, -7.22265625, -6.4921875], time:0.00305891ms
+   out_f16_th: [-6.69140625, -7.2265625, -6.4921875], time:0.00872254ms
 --------------------------------------------------------------------------------
 ```
diff --git a/kernels/hgemv/hgemv.cu b/kernels/hgemv/hgemv.cu
index 18791ac6..08b7e1cd 100644
--- a/kernels/hgemv/hgemv.cu
+++ b/kernels/hgemv/hgemv.cu
@@ -201,8 +201,15 @@ void hgemv_k16_f16(torch::Tensor a, torch::Tensor x, torch::Tensor y) {
                         reinterpret_cast<half *>(y.data_ptr()), M, K);
 }
 
+extern void hgemv_f16_cute(torch::Tensor, torch::Tensor, torch::Tensor);
+extern void hgemv_f16x8_cute(torch::Tensor, torch::Tensor, torch::Tensor);
+extern void hgemv_tensor_core_cute(torch::Tensor, torch::Tensor, torch::Tensor);
+
 PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
   TORCH_BINDING_COMMON_EXTENSION(hgemv_k32_f16)
   TORCH_BINDING_COMMON_EXTENSION(hgemv_k128_f16x4)
   TORCH_BINDING_COMMON_EXTENSION(hgemv_k16_f16)
+  TORCH_BINDING_COMMON_EXTENSION(hgemv_f16_cute)
+  TORCH_BINDING_COMMON_EXTENSION(hgemv_f16x8_cute)
+  TORCH_BINDING_COMMON_EXTENSION(hgemv_tensor_core_cute)
 }
diff --git a/kernels/hgemv/hgemv.py b/kernels/hgemv/hgemv.py
index 33f918d4..087851ad 100644
--- a/kernels/hgemv/hgemv.py
+++ b/kernels/hgemv/hgemv.py
@@ -1,3 +1,4 @@
+import os
 import time
 from functools import partial
 from typing import Optional
@@ -7,10 +8,14 @@
 
 torch.set_grad_enabled(False)
 
+CUTLASS_REPO_PATH = os.environ.get(
+    "CUTLASS_REPO_PATH", os.path.expanduser("../../third-party/cutlass")
+)
+
 # Load the CUDA kernel as a python module
 lib = load(
     name="hgemv_lib",
-    sources=["hgemv.cu"],
+    sources=["hgemv.cu", "hgemv_cute.cu"],
     extra_cuda_cflags=[
         "-O3",
         "-U__CUDA_NO_HALF_OPERATORS__",
@@ -22,6 +27,7 @@
         "--use_fast_math",
     ],
     extra_cflags=["-std=c++17"],
+    extra_include_paths=[os.path.join(CUTLASS_REPO_PATH, "include")],
 )
 
 
@@ -73,6 +79,9 @@ def run_benchmark(
 c = torch.randn((M, N)).cuda().half().contiguous()
 run_benchmark(lib.hgemv_k32_f16, a, b, "k32f16", c)
 run_benchmark(lib.hgemv_k128_f16x4, a, b, "k128f16x4", c)
+run_benchmark(lib.hgemv_f16_cute, a, b, "hgemv_f16_cute", c)
+run_benchmark(lib.hgemv_f16x8_cute, a, b, "hgemv_f16x8_cute", c)
+run_benchmark(lib.hgemv_tensor_core_cute, a, b, "hgemv_tensor_core_cute", c)
 run_benchmark(partial(torch.matmul, out=c), a, b, "f16_th")
 print("-" * 80)
 
@@ -81,5 +90,8 @@ def run_benchmark(
 b = torch.randn((K, N)).cuda().half().contiguous()
 c = torch.randn((M, N)).cuda().half().contiguous()
 run_benchmark(lib.hgemv_k16_f16, a, b, "k16f16", c)
+run_benchmark(lib.hgemv_f16_cute, a, b, "hgemv_f16_cute", c)
+run_benchmark(lib.hgemv_f16x8_cute, a, b, "hgemv_f16x8_cute", c)
+run_benchmark(lib.hgemv_tensor_core_cute, a, b, "hgemv_tensor_core_cute", c)
 run_benchmark(partial(torch.matmul, out=c), a, b, "f16_th")
 print("-" * 80)
diff --git a/kernels/hgemv/hgemv_cute.cu b/kernels/hgemv/hgemv_cute.cu
new file mode 100644
index 00000000..8ad9761d
--- /dev/null
+++ b/kernels/hgemv/hgemv_cute.cu
@@ -0,0 +1,408 @@
+﻿#include <cublas_v2.h>
+#include <cuda.h> // NOLINT
+
+#include <cute/layout.hpp>
+#include <cute/tensor.hpp>
+#include <stdlib.h>
+#include <torch/extension.h>
+
+using namespace cute;
+
+template <const int kWarpSize = 32>
+__device__ __forceinline__ half warp_reduce_sum_f16(half val) {
+#pragma unroll
+  for (int mask = kWarpSize >> 1; mask >= 1; mask >>= 1) {
+    val += __shfl_xor_sync(0xffffffff, val, mask);
+  }
+  return val;
+}
+
+template <typename T_, int NWarpPerBlock_> struct HgemvConfig {
+  using T = T_;
+  static constexpr int NWarpPerBlock = NWarpPerBlock_;
+  static constexpr int NumThreads = NWarpPerBlock * 32;
+
+  static constexpr int BlockM = 16 * NWarpPerBlock;
+  static constexpr int BlockN = 8;
+  static constexpr int BlockK = 16;
+
+  using MMA_Atom = MMA_Atom<SM80_16x8x16_F16F16F16F16_TN>;
+  using TiledMMA = decltype(make_tiled_mma(
+      MMA_Atom{},
+      make_layout(Shape<Int<NWarpPerBlock>, _1, _1>{}, GenColMajor{})));
+
+  static_assert(size(TiledMMA{}) == NumThreads && size(TiledMMA{}) <= 1024,
+                "NumThreads must be less than or equal 1024");
+};
+
+template <typename TiledCopy, int BlockM, int BlockK, int WARP_SIZE = 32>
+__global__ void hgemv_f16_cute_kernel(half *Aptr, half *Bptr, half *Cptr,
+                                      const int M, const int K) {
+  using namespace cute;
+
+  int thrid = threadIdx.x + threadIdx.y * blockDim.x;
+  int blockid = blockIdx.x;
+
+  int laneid = threadIdx.x % WARP_SIZE;
+  int warpid = threadIdx.y;
+
+  auto A = make_tensor(make_gmem_ptr(Aptr),
+                       make_layout(make_shape(M, K), make_stride(K, Int<1>{})));
+  auto B = make_tensor(make_gmem_ptr(Bptr),
+                       make_layout(make_shape(M, K), make_stride(0, Int<1>{})));
+  auto C = make_tensor(make_gmem_ptr(Cptr),
+                       make_layout(make_shape(M, 1), make_stride(Int<1>{}, 0)));
+
+  auto ABPre = make_identity_tensor(shape(A));
+  auto CPre = make_identity_tensor(shape(C));
+
+  auto gA = local_tile(A, make_shape(Int<BlockM>{}, Int<BlockK>{}),
+                       make_coord(blockid, _));
+  auto gB = local_tile(B, make_shape(Int<BlockM>{}, Int<BlockK>{}),
+                       make_coord(blockid, _));
+  auto gC = local_tile(C, make_shape(Int<BlockM>{}, Int<1>{}),
+                       make_coord(blockid, 0));
+
+  auto gABPre = local_tile(ABPre, make_shape(Int<BlockM>{}, Int<BlockK>{}),
+                           make_coord(blockid, _));
+  auto gCPre = local_tile(CPre, make_shape(Int<BlockM>{}, Int<1>{}),
+                          make_coord(blockid, _));
+
+  TiledCopy tiled_copy;
+  auto thr_copy = tiled_copy.get_slice(thrid);
+
+  auto tAgA = thr_copy.partition_S(gA);
+  auto tBgB = thr_copy.partition_S(gB);
+
+  auto rABPre = thr_copy.partition_S(gABPre);
+
+  int num_tile_k = size<2>(gA);
+
+  auto tArA = make_tensor_like(tAgA(_, _, _, 0));
+  auto tBrB = make_tensor_like(tBgB(_, _, _, 0));
+
+  auto sum = make_tensor_like(gC(0, _));
+  clear(sum);
+
+#pragma unroll
+  for (int num_iter_k = 0; num_iter_k < num_tile_k; num_iter_k++) {
+    auto pre_ = rABPre(_, _, _, num_iter_k);
+    auto pred = [&](auto... coords) {
+      return cute::elem_less(pre_(0), shape(A));
+    };
+
+    clear(tArA);
+    copy_if(tiled_copy, pred, tAgA(_, _, _, num_iter_k), tArA);
+    clear(tBrB);
+    copy_if(tiled_copy, pred, tBgB(_, _, _, num_iter_k), tBrB);
+
+    sum(0) += tArA(0) * tBrB(0);
+  }
+
+  sum(0) = warp_reduce_sum_f16<WARP_SIZE>(sum(0));
+
+  auto stord_pred = [&](auto... coords) {
+    return cute::elem_less(gCPre(warpid), shape(C)) && laneid == 0;
+  };
+  copy_if(stord_pred, sum, gC(warpid, _));
+}
+
+template <typename TiledCopy, int BlockM, int BlockK, int NumElemPerThread,
+          int WARP_SIZE = 32>
+__global__ void hgemv_f16x8_cute_kernel(half *Aptr, half *Bptr, half *Cptr,
+                                        const int M, const int K) {
+  using namespace cute;
+
+  int thrid = threadIdx.x + threadIdx.y * blockDim.x;
+  int blockid = blockIdx.x;
+
+  int laneid = threadIdx.x % WARP_SIZE;
+  int warpid = threadIdx.y;
+
+  auto A = make_tensor(make_gmem_ptr(Aptr),
+                       make_layout(make_shape(M, K), make_stride(K, Int<1>{})));
+  auto B = make_tensor(make_gmem_ptr(Bptr),
+                       make_layout(make_shape(M, K), make_stride(0, Int<1>{})));
+  auto C = make_tensor(make_gmem_ptr(Cptr),
+                       make_layout(make_shape(M, 1), make_stride(Int<1>{}, 0)));
+
+  auto ABPre = make_identity_tensor(shape(A));
+  auto CPre = make_identity_tensor(shape(C));
+
+  auto gA = local_tile(A, make_shape(Int<BlockM>{}, Int<BlockK>{}),
+                       make_coord(blockid, _));
+  auto gB = local_tile(B, make_shape(Int<BlockM>{}, Int<BlockK>{}),
+                       make_coord(blockid, _));
+  auto gC = local_tile(C, make_shape(Int<BlockM>{}, Int<1>{}),
+                       make_coord(blockid, 0));
+
+  auto gABPre = local_tile(ABPre, make_shape(Int<BlockM>{}, Int<BlockK>{}),
+                           make_coord(blockid, _));
+  auto gCPre = local_tile(CPre, make_shape(Int<BlockM>{}, Int<1>{}),
+                          make_coord(blockid, _));
+
+  TiledCopy tiled_copy;
+  auto thr_copy = tiled_copy.get_slice(thrid);
+
+  auto tAgA = thr_copy.partition_S(gA);
+  auto tBgB = thr_copy.partition_S(gB);
+  auto rABPre = thr_copy.partition_S(gABPre);
+
+  int num_tile_k = size<2>(gA);
+
+  auto tArA = make_tensor_like(tAgA(_, _, _, 0));
+  auto tBrB = make_tensor_like(tBgB(_, _, _, 0));
+
+  auto sum = make_tensor_like(gC(0, _));
+  clear(sum);
+
+#pragma unroll
+  for (int iter_k = 0; iter_k < num_tile_k; iter_k++) {
+    auto pre_ = rABPre(_, _, _, iter_k);
+    auto pred = [&](auto... coords) {
+      return cute::elem_less(pre_(NumElemPerThread - 1), shape(A));
+    };
+
+    clear(tArA);
+    copy_if(tiled_copy, pred, tAgA(_, _, _, iter_k), tArA);
+    clear(tBrB);
+    copy_if(tiled_copy, pred, tBgB(_, _, _, iter_k), tBrB);
+
+    auto tArA_half2 = recast<half2>(tArA);
+    auto tBrB_half2 = recast<half2>(tBrB);
+    auto sum_half2 = make_tensor<half2>(make_shape(Int<1>{}));
+
+#pragma unroll
+    for (int iter_elem = 0; iter_elem < size(tArA_half2); iter_elem++) {
+      sum_half2(0) =
+          tArA_half2(iter_elem) * tBrB_half2(iter_elem) + sum_half2(0);
+    }
+
+    sum(0) += sum_half2(0).x + sum_half2(0).y;
+  }
+
+  sum(0) = warp_reduce_sum_f16<WARP_SIZE>(sum(0));
+
+  auto stord_pred = [&](auto... coords) {
+    return cute::elem_less(gCPre(warpid), shape(C)) && laneid == 0;
+  };
+  copy_if(stord_pred, sum, gC(warpid, _));
+}
+
+// using tensor core
+template <typename HgemvConfig_>
+__global__ void hgemv_tensor_core_cute_kernel(typename HgemvConfig_::T *Aptr,
+                                              typename HgemvConfig_::T *Bptr,
+                                              typename HgemvConfig_::T *Cptr,
+                                              const int M, const int K) {
+  using namespace cute;
+
+  using T = typename HgemvConfig_::T;
+  using TiledMMA = typename HgemvConfig_::TiledMMA;
+  constexpr int BlockM = HgemvConfig_::BlockM;
+  constexpr int BlockN = HgemvConfig_::BlockN;
+  constexpr int BlockK = HgemvConfig_::BlockK;
+
+  int thrid = threadIdx.x;
+  int blockid = blockIdx.x;
+
+  int warpid = threadIdx.x / 32;
+  int laneid = threadIdx.x % 32;
+
+  auto A = make_tensor(make_gmem_ptr(Aptr),
+                       make_layout(make_shape(M, K), make_stride(K, Int<1>{})));
+  auto B = make_tensor(make_gmem_ptr(Bptr),
+                       make_layout(make_shape(M, K), make_stride(0, Int<1>{})));
+  auto C = make_tensor(make_gmem_ptr(Cptr),
+                       make_layout(make_shape(M, 1), make_stride(Int<1>{}, 0)));
+
+  auto ABPre = make_identity_tensor(shape(A));
+  auto CPre = make_identity_tensor(shape(C));
+
+  auto gA = local_tile(A, make_shape(Int<BlockM>{}, Int<BlockK>{}),
+                       make_coord(blockid, _));
+  auto gB = local_tile(B, make_shape(Int<BlockN>{}, Int<BlockK>{}),
+                       make_coord(blockid, _));
+  auto gC = local_tile(C, make_shape(Int<BlockM>{}, Int<1>{}),
+                       make_coord(blockid, 0));
+
+  auto gABPre = local_tile(ABPre, make_shape(Int<BlockM>{}, Int<BlockK>{}),
+                           make_coord(blockid, _));
+  auto gCPre = local_tile(CPre, make_shape(Int<BlockM>{}, Int<1>{}),
+                          make_coord(blockid, _));
+
+  TiledMMA tiled_mma;
+  auto thr_mma = tiled_mma.get_slice(thrid);
+  auto tAgA = thr_mma.partition_A(gA);
+  auto tBgB = thr_mma.partition_B(gB);
+
+  auto rAPre = thr_mma.partition_A(gABPre);
+  auto rBPre = thr_mma.partition_B(gABPre);
+
+  auto tArA = make_tensor_like(tAgA(_, _, _, 0));
+  auto tBrB = make_tensor_like(tBgB(_, _, _, 0));
+
+  auto tCrC =
+      partition_fragment_C(tiled_mma, Shape<Int<BlockM>, Int<BlockN>>{});
+
+  clear(tCrC);
+
+  int num_tile_k = size<2>(gA);
+#pragma unroll
+  for (int itile = 0; itile < num_tile_k; itile++) {
+    auto pre_A = rAPre(_, _, _, itile);
+    auto pre_B = rBPre(_, _, _, itile);
+    auto pred_A = [&](auto... coords) {
+      return cute::elem_less(pre_A(coords...), shape(A));
+    };
+    auto pred_B = [&](auto... coords) {
+      return cute::elem_less(pre_B(coords...), shape(A));
+    };
+
+    clear(tArA);
+    copy_if(pred_A, tAgA(_, _, _, itile), tArA);
+    clear(tBrB);
+    copy_if(pred_B, tBgB(_, _, _, itile), tBrB);
+
+    gemm(tiled_mma, tArA, tBrB, tCrC);
+  }
+
+  int elem_index1 = warpid * 16 + laneid / 4;
+  int elem_index2 = warpid * 16 + laneid / 4 + 8;
+
+  auto sum = make_tensor_like(gC(0, _));
+  sum(0) = tCrC(0);
+  auto elem_pred1 = [&](auto... coords) {
+    return (laneid % 4 == 0) && cute::elem_less(gCPre(elem_index1), shape(C));
+  };
+  copy_if(elem_pred1, sum, gC(elem_index1, _));
+
+  sum(0) = tCrC(2);
+  auto elem_pred2 = [&](auto... coords) {
+    return (laneid % 4 == 0) && cute::elem_less(gCPre(elem_index2), shape(C));
+  };
+  copy_if(elem_pred2, sum, gC(elem_index2, _));
+}
+
+// --------------------- PyTorch bindings for custom kernel
+// -----------------------
+#define STRINGFY(str) #str
+#define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
+  m.def(STRINGFY(func), &func, STRINGFY(func));
+
+#define CHECK_TORCH_TENSOR_DTYPE(T, th_type)                                   \
+  if (((T).options().dtype() != (th_type))) {                                  \
+    std::cout << "Tensor Info:" << (T).options() << std::endl;                 \
+    throw std::runtime_error("values must be " #th_type);                      \
+  }
+
+#define CHECK_TORCH_TENSOR_SHAPE(T, S0, S1)                                    \
+  if (((T).size(0) != (S0)) || ((T).size(1) != (S1))) {                        \
+    throw std::runtime_error("Tensor size mismatch!");                         \
+  }
+
+#define ASSERT_K_IS_MULTIBLE_OF(V)                                             \
+  if (K % (V) != 0) {                                                          \
+    throw std::runtime_error("K must be multiple of " #V);                     \
+  }
+
+#define ASSERT_K_IS_EQUAL_OF(V)                                                \
+  if (K != (V)) {                                                              \
+    throw std::runtime_error("K must be " #V);                                 \
+  }
+
+void hgemv_f16_cute(torch::Tensor A, torch::Tensor B, torch::Tensor C) {
+
+  CHECK_TORCH_TENSOR_DTYPE(A, torch::kHalf)
+  CHECK_TORCH_TENSOR_DTYPE(B, torch::kHalf)
+  CHECK_TORCH_TENSOR_DTYPE(C, torch::kHalf)
+  const int M = A.size(0);
+  const int K = A.size(1);
+  CHECK_TORCH_TENSOR_SHAPE(A, M, K)
+  CHECK_TORCH_TENSOR_SHAPE(B, K, 1)
+  CHECK_TORCH_TENSOR_SHAPE(C, M, 1)
+  // ASSERT_K_IS_MULTIBLE_OF(8)
+
+  constexpr int NumThreadPerRow = 32;
+  constexpr int NumThreadPerBlock = 128;
+  constexpr int NumRowPerBlcok = NumThreadPerBlock / 32;
+
+  using LoadType = uint16_t;
+
+  constexpr int NumElemPerThread = sizeof(LoadType) / sizeof(half);
+
+  using CopyAtom = Copy_Atom<UniversalCopy<LoadType>, half>;
+  using TiledCopy = decltype(make_tiled_copy(
+      CopyAtom{},
+      make_layout(Shape<Int<NumRowPerBlcok>, Int<NumThreadPerRow>>{},
+                  GenRowMajor{}),
+      make_layout(Shape<_1, Int<NumElemPerThread>>{}, GenRowMajor{})));
+
+  dim3 blcok(NumThreadPerRow, NumRowPerBlcok);
+  dim3 grid(ceil_div(M, NumRowPerBlcok));
+
+  hgemv_f16_cute_kernel<TiledCopy, NumRowPerBlcok,
+                        NumThreadPerRow * NumElemPerThread>
+      <<<grid, blcok>>>(reinterpret_cast<half *>(A.data_ptr()),
+                        reinterpret_cast<half *>(B.data_ptr()),
+                        reinterpret_cast<half *>(C.data_ptr()), M, K);
+}
+
+void hgemv_f16x8_cute(torch::Tensor A, torch::Tensor B, torch::Tensor C) {
+  CHECK_TORCH_TENSOR_DTYPE(A, torch::kHalf)
+  CHECK_TORCH_TENSOR_DTYPE(B, torch::kHalf)
+  CHECK_TORCH_TENSOR_DTYPE(C, torch::kHalf)
+  const int M = A.size(0);
+  const int K = A.size(1);
+  CHECK_TORCH_TENSOR_SHAPE(A, M, K)
+  CHECK_TORCH_TENSOR_SHAPE(B, K, 1)
+  CHECK_TORCH_TENSOR_SHAPE(C, M, 1)
+  ASSERT_K_IS_MULTIBLE_OF(8)
+
+  constexpr int NumThreadPerRow = 32;
+  constexpr int NumThreadPerBlock = 128;
+  constexpr int NumRowPerBlcok = NumThreadPerBlock / 32;
+
+  using LoadType = uint128_t;
+
+  constexpr int NumElemPerThread = sizeof(LoadType) / sizeof(half);
+
+  using CopyAtom = Copy_Atom<UniversalCopy<LoadType>, half>;
+  using TiledCopy = decltype(make_tiled_copy(
+      CopyAtom{},
+      make_layout(Shape<Int<NumRowPerBlcok>, Int<NumThreadPerRow>>{},
+                  GenRowMajor{}),
+      make_layout(Shape<_1, Int<NumElemPerThread>>{}, GenRowMajor{})));
+
+  dim3 blcok(NumThreadPerRow, NumRowPerBlcok);
+  dim3 grid(ceil_div(M, NumRowPerBlcok));
+
+  hgemv_f16x8_cute_kernel<TiledCopy, NumRowPerBlcok,
+                          NumThreadPerRow * NumElemPerThread, NumElemPerThread>
+      <<<grid, blcok>>>(reinterpret_cast<half *>(A.data_ptr()),
+                        reinterpret_cast<half *>(B.data_ptr()),
+                        reinterpret_cast<half *>(C.data_ptr()), M, K);
+}
+
+void hgemv_tensor_core_cute(torch::Tensor A, torch::Tensor B, torch::Tensor C) {
+
+  CHECK_TORCH_TENSOR_DTYPE(A, torch::kHalf)
+  CHECK_TORCH_TENSOR_DTYPE(B, torch::kHalf)
+  CHECK_TORCH_TENSOR_DTYPE(C, torch::kHalf)
+  const int M = A.size(0);
+  const int K = A.size(1);
+  CHECK_TORCH_TENSOR_SHAPE(A, M, K)
+  CHECK_TORCH_TENSOR_SHAPE(B, K, 1)
+  CHECK_TORCH_TENSOR_SHAPE(C, M, 1)
+  // ASSERT_K_IS_MULTIBLE_OF(8)
+
+  using config = HgemvConfig<half, 4>;
+
+  dim3 blcok(size(config::NumThreads));
+  dim3 grid(ceil_div(M, config::BlockM));
+
+  hgemv_tensor_core_cute_kernel<config>
+      <<<grid, blcok>>>(reinterpret_cast<half *>(A.data_ptr()),
+                        reinterpret_cast<half *>(B.data_ptr()),
+                        reinterpret_cast<half *>(C.data_ptr()), M, K);
+}

From 8e04cde24fc4f63e036bd638dae7d53821d90745 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Sat, 7 Jun 2025 09:49:05 +0800
Subject: [PATCH 09/57] Update README.md (#333)

---
 README.md | 14 +++++++-------
 1 file changed, 7 insertions(+), 7 deletions(-)

diff --git a/README.md b/README.md
index 6a31742b..437381ed 100644
--- a/README.md
+++ b/README.md
@@ -57,13 +57,6 @@
 ## 📖 Contents
 <div id="contents"></div>
 <!---
-- [📖 HGEMM-MMA 🎉🎉](#HGEMM-bench)
-- [📖 FlashAttention-MMA 🎉🎉](#fa-mma-bench)
-- [📖 200+ CUDA Kernels 🔥🔥](#cuda-kernel)
-- [📖 100+ 高性能计算文章 💡💡](#my-blogs-part-1)
-- [📖 How to Contribute 👀👇](#contribute)
---->
-
 - [📖 HGEMM-MMA 🎉🎉](#HGEMM-bench)
   - [📚 CUDA/Tensor Cores](#HGEMM-bench)
   - [📚 Tile Block(Br, Bc)](#HGEMM-bench)
@@ -75,6 +68,13 @@
   - [📚 Register Double Buffers](#HGEMM-bench)
   - [📚 Collective Store(Shfl)](#HGEMM-bench)
   - [📚 Layout NN/TN](#HGEMM-bench)
+- [📖 FlashAttention-MMA 🎉🎉](#fa-mma-bench)
+- [📖 200+ CUDA Kernels 🔥🔥](#cuda-kernel)
+- [📖 100+ 高性能计算文章 💡💡](#my-blogs-part-1)
+- [📖 How to Contribute 👀👇](#contribute)
+--->
+
+- [📖 HGEMM-MMA 🎉🎉](#HGEMM-bench)
 - [📖 FlashAttention-MMA 🎉🎉](#fa-mma-bench)
   - [📚 Split KV (Basic, FA-1)](#mma-split-kv)
   - [📚 Split Q (Faster, FA-2)](#mma-split-q)

From 491efab1463d4ed5a6f29b7110cb5b5ffc48cb9c Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Sat, 7 Jun 2025 10:12:55 +0800
Subject: [PATCH 10/57] Update README.md

---
 README.md | 16 ++++++++--------
 1 file changed, 8 insertions(+), 8 deletions(-)

diff --git a/README.md b/README.md
index 437381ed..b0166668 100644
--- a/README.md
+++ b/README.md
@@ -71,6 +71,14 @@
 - [📖 FlashAttention-MMA 🎉🎉](#fa-mma-bench)
 - [📖 200+ CUDA Kernels 🔥🔥](#cuda-kernel)
 - [📖 100+ 高性能计算文章 💡💡](#my-blogs-part-1)
+  - [📚 大模型推理优化原理](#my-blogs-part-1)
+  - [📚 大模型分布式训推原理](#my-blogs-part-1)
+  - [📚 CV/C++/模型部署优化](#my-blogs-part-1)
+  - [📚 CUDA优化入门与实践](#other-blogs)
+  - [📚 Tensor Cores入门教程](#other-blogs)
+  - [📚 CuTe系列详解与实践](#other-blogs)
+  - [📚 GPU指令集架构精解](#other-blogs)
+  - [📚 GPU通信架构精解](#other-blogs)
 - [📖 How to Contribute 👀👇](#contribute)
 --->
 
@@ -91,14 +99,6 @@
   - [📚 Triton ⭐⭐⭐️](#triton-kernel)
   - [📚 CUTLASS ⭐⭐⭐️](#cutlass-kernel)
 - [📖 100+ 高性能计算文章 💡💡](#my-blogs-part-1)
-  - [📚 大模型推理优化原理](#my-blogs-part-1)
-  - [📚 大模型分布式训推原理](#my-blogs-part-1)
-  - [📚 CV/C++/模型部署优化](#my-blogs-part-1)
-  - [📚 CUDA优化入门与实践](#other-blogs)
-  - [📚 Tensor Cores入门教程](#other-blogs)
-  - [📚 CuTe系列详解与实践](#other-blogs)
-  - [📚 GPU指令集架构精解](#other-blogs)
-  - [📚 GPU通信架构精解](#other-blogs)
 - [📖 How to Contribute 👀👇](#contribute)
 
 

From 38887759b4e9529fec7696ba4498664ff7914284 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Sat, 7 Jun 2025 10:16:37 +0800
Subject: [PATCH 11/57] Update README.md

---
 README.md | 156 +++++++++++++++++++++++++++---------------------------
 1 file changed, 78 insertions(+), 78 deletions(-)

diff --git a/README.md b/README.md
index b0166668..51765337 100644
--- a/README.md
+++ b/README.md
@@ -228,12 +228,12 @@ flash_attn_mma_stages_split_q_tiling_qkv_kernel(half* Q, half* K, half* V, half*
 ## ©️Citations🎉🎉
 
 ```BibTeX
-@misc{LeetCUDA@2024,
+@misc{LeetCUDA@2025,
   title={LeetCUDA: A Modern CUDA Learn Notes with PyTorch for Beginners},
   url={https://github.com/xlite-dev/LeetCUDA},
   note={Open-source software available at https://github.com/xlite-dev/LeetCUDA},
-  author={xlite-dev etc},
-  year={2024}
+  author={DefTruth etc},
+  year={2025}
 }
 ```
 
@@ -509,81 +509,81 @@ The kernels listed here will guide you through a step-by-step progression, rangi
 
 |📖 类型-标题|📖 作者| 📖 推荐 |
 |:---|:---|:---|
-| [[Triton编程][基础]📖Triton极简入门: Triton Vector Add](https://zhuanlan.zhihu.com/p/1902778199261291694)|@xlite-dev|⭐️⭐️⭐|
-| [[Triton编程][基础]📖Triton Fused Softmax Kernel详解: 从Python源码到PTX分析](https://zhuanlan.zhihu.com/p/1899562146477609112)|@xlite-dev|⭐️⭐️⭐|
-| [[Triton编程][基础]📖vLLM Triton Merge Attention States Kernel详解](https://zhuanlan.zhihu.com/p/1904937907703243110)|@xlite-dev|⭐️⭐️⭐|
-| [[Triton编程][进阶]📖vLLM Prefix Prefill Triton Kernel图解](https://zhuanlan.zhihu.com/p/695799736)|@xlite-dev|⭐️⭐️⭐️|
-| [[张量/序列并行]📖序列并行: BPT、Ring-Attention及Striped-Attention笔记](https://zhuanlan.zhihu.com/p/6456708235)|@xlite-dev|⭐️⭐️⭐|
-| [[vLLM实践][算子]📖vLLM算子开发流程：”保姆级“详细记录](https://zhuanlan.zhihu.com/p/1892966682634473987)|@xlite-dev|⭐️⭐️⭐|
-| [[vLLM实践][万字]📖vLLM + DeepSeek-R1 671B 多机部署及修Bug笔记](https://zhuanlan.zhihu.com/p/29950052712)|@xlite-dev|⭐️⭐️⭐|
-| [[Attention优化]📖FFPA(Split-D): FA2无限HeadDim扩展，2x↑🎉 vs SDPA EA](https://zhuanlan.zhihu.com/p/13975660308)|@xlite-dev|⭐️⭐️⭐️|
-| [[CUDA基础][开篇]📖LeetCUDA: v3.0 大升级-面试刷题不迷路](https://zhuanlan.zhihu.com/p/19862356369)|@xlite-dev|⭐️⭐️⭐⭐️|
-| [[分布式训推][张量/序列并行]📖图解DeepSpeed-Ulysses&Megatron-LM TP/SP](https://zhuanlan.zhihu.com/p/5750410146)|@xlite-dev|⭐️⭐️|
-| [[VLM推理优化][InternVL系列]📖InternLM2/.../InternVL1.5系列笔记: 核心点解析](https://zhuanlan.zhihu.com/p/702481058)|@xlite-dev|⭐️⭐️|
-| [[LLM推理优化][TensorRT-LLM][5w字]📖TensorRT-LLM部署调优-指北](https://zhuanlan.zhihu.com/p/699333691)|@xlite-dev|⭐️⭐️⭐️|
-| [[LLM推理优化][KV Cache优化]📖GQA/YOCO/CLA/MLKV: 层内和层间KV Cache共享](https://zhuanlan.zhihu.com/p/697311739)|@xlite-dev|⭐️⭐️|
-| [[LLM推理优化][Prefill优化][万字]📖图解vLLM Automatic Prefix Caching: TTFT优化](https://zhuanlan.zhihu.com/p/693556044)|@xlite-dev|⭐️⭐️⭐️|
-| [[LLM推理优化][Attention优化]📖图解:从Online-Softmax到FlashAttention V1/V2/V3](https://zhuanlan.zhihu.com/p/668888063)|@xlite-dev|⭐️⭐️⭐️|
-| [[LLM推理优化][Decoding优化]📖原理&图解FlashDecoding/FlashDecoding++](https://zhuanlan.zhihu.com/p/696075602)|@xlite-dev|⭐️⭐️|
-| [[VLM推理优化][LLaVA系列]📖CLIP/LLaVA/LLaVA1.5/VILA笔记: 核心点解析](https://zhuanlan.zhihu.com/p/683137074)|@xlite-dev|⭐️⭐️|
-| [[LLM推理优化][Attention优化][万字]📖TensorRT MHA/Myelin vs FlashAttention-2](https://zhuanlan.zhihu.com/p/678873216)|@xlite-dev|⭐️⭐️⭐️|
-| [[LLM推理优化][PTX汇编]📖CUDA 12 PTX汇编: PRMT指令详解-通用模式](https://zhuanlan.zhihu.com/p/660630414)|@xlite-dev|⭐️|
-| [[LLM推理优化][PTX汇编]📖CUDA 12 PTX汇编: LOP3指令详解](https://zhuanlan.zhihu.com/p/659741469)|@xlite-dev|⭐️|
-| [[LLM推理优化][CUDA][3w字]📖高频面试题汇总-大模型手撕CUDA](https://zhuanlan.zhihu.com/p/678903537)|@xlite-dev|⭐️⭐️⭐️|
-| [[LLM推理优化][Weight Only]📖WINT8/4-(00): 通俗易懂讲解-快速反量化算法](https://zhuanlan.zhihu.com/p/657072856)|@xlite-dev|⭐️⭐️|
-| [[LLM推理优化][Weight Only]📖WINT8/4-(01): PRMT指令详解及FT源码解析](https://zhuanlan.zhihu.com/p/657070837)|@xlite-dev|⭐️⭐️|
-| [[LLM推理优化][Weight Only]📖WINT8/4-(02): 快速反量化之INT8转BF16](https://zhuanlan.zhihu.com/p/657073159)|@xlite-dev|⭐️⭐️|
-| [[LLM推理优化][Weight Only]📖WINT8/4-(03): LOP3指令详解及INT4转FP16/BF16](https://zhuanlan.zhihu.com/p/657073857)|@xlite-dev|⭐️⭐️|
-| [[LLM推理优化][LLM Infra整理]📖100+篇: 大模型推理各方向新发展整理](https://zhuanlan.zhihu.com/p/693680304)|@xlite-dev|⭐️⭐️|
-| [[LLM推理优化][LLM Infra整理]📖30+篇: LLM推理论文集-500页PDF](https://zhuanlan.zhihu.com/p/669777159)|@xlite-dev|⭐️⭐️|
-| [[LLM推理优化][LLM Infra整理]📖FlashDecoding++: 比FlashDecoding还要快！](https://zhuanlan.zhihu.com/p/665022589)|@xlite-dev|⭐️|
-| [[LLM推理优化][LLM Infra整理]📖TensorRT-LLM开源，TensorRT 9.1也来了](https://zhuanlan.zhihu.com/p/662361469)|@xlite-dev|⭐️|
-| [[LLM推理优化][LLM Infra整理]📖20+篇: LLM推理论文集-300页PDF](https://zhuanlan.zhihu.com/p/658091768)|@xlite-dev|⭐️⭐️|
-| [[LLM推理优化][LLM Infra整理]📖PagedAttention论文新鲜出炉](https://zhuanlan.zhihu.com/p/617015570)|@xlite-dev|⭐️|
-| [[推理部署][CV/NLP]📖FastDeploy三行代码搞定150+ CV、NLP模型部署](https://zhuanlan.zhihu.com/p/581326442)|@xlite-dev|⭐️|
-| [[推理部署][CV]📖如何在lite.ai.toolkit(3.6k+ stars)中增加您的模型？](https://zhuanlan.zhihu.com/p/523876625)|@xlite-dev|⭐️⭐️|
-| [[推理部署][CV]📖美团 YOLOv6 ORT/MNN/TNN/NCNN C++推理部署](https://zhuanlan.zhihu.com/p/533643238)|@xlite-dev|⭐️⭐️|
-| [[推理部署][ONNX]📖ONNX推理加速技术文档-杂记](https://zhuanlan.zhihu.com/p/524023964)|@xlite-dev|⭐️|
-| [[推理部署][TensorFlow]📖Mac源码编译TensorFlow C++指北](https://zhuanlan.zhihu.com/p/524013615)|@xlite-dev|⭐️|
-| [[推理部署][CV]📖1Mb!头部姿态估计: FSANet，一个小而美的模型(C++)](https://zhuanlan.zhihu.com/p/447364201)|@xlite-dev|⭐️|
-| [[推理部署][CV]📖opencv+ffmpeg编译打包全解指南](https://zhuanlan.zhihu.com/p/472115312)|@xlite-dev|⭐️⭐️|
-| [[推理部署][CV]📖RobustVideoMatting视频抠图静态ONNX模型转换](https://zhuanlan.zhihu.com/p/459088407)|@xlite-dev|⭐️|
-| [[推理部署][CV]📖190Kb!SSRNet年龄检测详细解读（含C++工程）](https://zhuanlan.zhihu.com/p/462762797)|@xlite-dev|⭐️|
-| [[推理部署][CV]📖MGMatting(CVPR2021)人像抠图C++应用记录](https://zhuanlan.zhihu.com/p/464732042)|@xlite-dev|⭐️|
-| [[推理部署][CV]📖超准确人脸检测(带关键点)YOLO5Face C++工程详细记录](https://zhuanlan.zhihu.com/p/461878005)|@xlite-dev|⭐️⭐️|
-| [[推理部署][ORT]📖解决: ONNXRuntime(Python) GPU 部署配置记录](https://zhuanlan.zhihu.com/p/457484536)|@xlite-dev|⭐️|
-| [[推理部署][CV]📖记录SCRFD(CVPR2021)人脸检测C++工程化(含docker镜像)](https://zhuanlan.zhihu.com/p/455165568)|@xlite-dev|⭐️⭐️|
-| [[推理部署][NCNN]📖野路子：记录一个解决onnx转ncnn时op不支持的trick](https://zhuanlan.zhihu.com/p/451446147)|@xlite-dev|⭐️|
-| [[推理部署][CV]📖升级版轻量级NanoDet-Plus MNN/TNN/NCNN/ORT C++工程记录](https://zhuanlan.zhihu.com/p/450586647)|@xlite-dev|⭐️⭐️|
-| [[推理部署][CV]📖超轻量级NanoDet MNN/TNN/NCNN/ORT C++工程记录](https://zhuanlan.zhihu.com/p/443419387)|@xlite-dev|⭐️|
-| [[推理部署][CV]📖详细记录MGMatting之MNN、TNN和ORT C++移植](https://zhuanlan.zhihu.com/p/442949027)|@xlite-dev|⭐️⭐️|
-| [[推理部署][CV]📖YOLOX NCNN/MNN/TNN/ONNXRuntime C++工程简记](https://zhuanlan.zhihu.com/p/447364122)|@xlite-dev|⭐️|
-| [[推理部署][TNN]📖手动修改YoloX的tnnproto记录-TNN](https://zhuanlan.zhihu.com/p/425668734)|@xlite-dev|⭐️|
-| [[推理部署][ORT]📖全网最详细 ONNXRuntime C++/Java/Python 资料！](https://zhuanlan.zhihu.com/p/414317269)|@xlite-dev|⭐️|
-| [[推理部署][CV]📖RobustVideoMatting: C++工程化记录-实现篇](https://zhuanlan.zhihu.com/p/413280488)|@xlite-dev|⭐️⭐️|
-| [[推理部署][CV]📖RobustVideoMatting: C++工程化记录-应用篇](https://zhuanlan.zhihu.com/p/412491918)|@xlite-dev|⭐️⭐️|
-| [[推理部署][ORT]📖ONNXRuntime C++ CMake 工程分析及编译](https://zhuanlan.zhihu.com/p/411887386)|@xlite-dev|⭐️⭐️|
-| [[推理部署][ORT]📖如何使用ORT C++ API处理NCHW和NHWC输入？](https://zhuanlan.zhihu.com/p/524230808)|@xlite-dev|⭐️|
-| [[推理部署][TNN]📖tnn-convert搭建简记-YOLOP转TNN](https://zhuanlan.zhihu.com/p/431418709)|@xlite-dev|⭐️|
-| [[推理部署][CV]📖YOLOP ONNXRuntime C++工程化记录](https://zhuanlan.zhihu.com/p/411651933)|@xlite-dev|⭐️⭐️|
-| [[推理部署][NCNN]📖超有用NCNN参考资料整理](https://zhuanlan.zhihu.com/p/449765328)|@xlite-dev|⭐️|
-| [[推理部署][MNN]📖超有用MNN参考资料整理](https://zhuanlan.zhihu.com/p/449761992)|@xlite-dev|⭐️|
-| [[推理部署][TNN]📖超有用TNN参考资料整理](https://zhuanlan.zhihu.com/p/449769615)|@xlite-dev|⭐️|
-| [[推理部署][ONNX]📖超有用ONNX参考资料整理](https://zhuanlan.zhihu.com/p/449773663)|@xlite-dev|⭐️|
-| [[推理部署][ONNX]📖超有用ONNX模型结构参考资料整理](https://zhuanlan.zhihu.com/p/449775926)|@xlite-dev|⭐️|
-| [[推理部署][OpenCV-DNN]📖超有用OpenCV-DNN参考资料整理](https://zhuanlan.zhihu.com/p/449778377)|@xlite-dev|⭐️|
-| [[推理部署][Tensorflow]📖超有用Tensorflow C++工程化知识点](https://zhuanlan.zhihu.com/p/449788027)|@xlite-dev|⭐️|
-| [[推理部署][模型转换]📖深度学习模型转换资料整理](https://zhuanlan.zhihu.com/p/449759361)|@xlite-dev|⭐️|
-| [[技术随笔][C++][CMake]📖超有用CMake参考资料整理](https://zhuanlan.zhihu.com/p/449779892)|@xlite-dev|⭐️⭐️|
-| [[技术随笔][C++][3W字]📖静态链接和静态库实践指北-原理篇](https://zhuanlan.zhihu.com/p/595527528)|@xlite-dev|⭐️⭐️⭐️|
-| [[技术随笔][C++]📖Mac下C++内存检查指北(Valgrind VS Asan)](https://zhuanlan.zhihu.com/p/508470880)|@xlite-dev|⭐️|
-| [[技术随笔][CV]📖torchlm: 人脸关键点检测库](https://zhuanlan.zhihu.com/p/467211561)|@xlite-dev|⭐️⭐️|
-| [[技术随笔][ML]📖《统计学习方法-李航: 笔记-从原理到实现-基于R》](https://zhuanlan.zhihu.com/p/684885595)|@xlite-dev|⭐️⭐️|
-| [[技术随笔][Git]📖如何优雅地git clone和git submodule？](https://zhuanlan.zhihu.com/p/639136221)|@xlite-dev|⭐️|
-| [[技术随笔][3D]📖人脸重建3D参考资料整理](https://zhuanlan.zhihu.com/p/524034741)|@xlite-dev|⭐️|
-| [[技术随笔][3D]📖BlendShapes参考资料整理](https://zhuanlan.zhihu.com/p/524036145)|@xlite-dev|⭐️|
-| [[技术随笔][3D]📖从源码安装Pytorch3D详细记录及学习资料](https://zhuanlan.zhihu.com/p/512347464)|@xlite-dev|⭐️|
-| [[技术随笔][ML]📖200页:《统计学习方法：李航》笔记 -从原理到实现](https://zhuanlan.zhihu.com/p/461520847)|@xlite-dev|⭐️⭐️|
+| [[Triton编程][基础]📖Triton极简入门: Triton Vector Add](https://zhuanlan.zhihu.com/p/1902778199261291694)|@DefTruth|⭐️⭐️⭐|
+| [[Triton编程][基础]📖Triton Fused Softmax Kernel详解: 从Python源码到PTX分析](https://zhuanlan.zhihu.com/p/1899562146477609112)|@DefTruth|⭐️⭐️⭐|
+| [[Triton编程][基础]📖vLLM Triton Merge Attention States Kernel详解](https://zhuanlan.zhihu.com/p/1904937907703243110)|@DefTruth|⭐️⭐️⭐|
+| [[Triton编程][进阶]📖vLLM Prefix Prefill Triton Kernel图解](https://zhuanlan.zhihu.com/p/695799736)|@DefTruth|⭐️⭐️⭐️|
+| [[张量/序列并行]📖序列并行: BPT、Ring-Attention及Striped-Attention笔记](https://zhuanlan.zhihu.com/p/6456708235)|@DefTruth|⭐️⭐️⭐|
+| [[vLLM实践][算子]📖vLLM算子开发流程：”保姆级“详细记录](https://zhuanlan.zhihu.com/p/1892966682634473987)|@DefTruth|⭐️⭐️⭐|
+| [[vLLM实践][万字]📖vLLM + DeepSeek-R1 671B 多机部署及修Bug笔记](https://zhuanlan.zhihu.com/p/29950052712)|@DefTruth|⭐️⭐️⭐|
+| [[Attention优化]📖FFPA(Split-D): FA2无限HeadDim扩展，2x↑🎉 vs SDPA EA](https://zhuanlan.zhihu.com/p/13975660308)|@DefTruth|⭐️⭐️⭐️|
+| [[CUDA基础][开篇]📖LeetCUDA: v3.0 大升级-面试刷题不迷路](https://zhuanlan.zhihu.com/p/19862356369)|@DefTruth|⭐️⭐️⭐⭐️|
+| [[分布式训推][张量/序列并行]📖图解DeepSpeed-Ulysses&Megatron-LM TP/SP](https://zhuanlan.zhihu.com/p/5750410146)|@DefTruth|⭐️⭐️|
+| [[VLM推理优化][InternVL系列]📖InternLM2/.../InternVL1.5系列笔记: 核心点解析](https://zhuanlan.zhihu.com/p/702481058)|@DefTruth|⭐️⭐️|
+| [[LLM推理优化][TensorRT-LLM][5w字]📖TensorRT-LLM部署调优-指北](https://zhuanlan.zhihu.com/p/699333691)|@DefTruth|⭐️⭐️⭐️|
+| [[LLM推理优化][KV Cache优化]📖GQA/YOCO/CLA/MLKV: 层内和层间KV Cache共享](https://zhuanlan.zhihu.com/p/697311739)|@DefTruth|⭐️⭐️|
+| [[LLM推理优化][Prefill优化][万字]📖图解vLLM Automatic Prefix Caching: TTFT优化](https://zhuanlan.zhihu.com/p/693556044)|@DefTruth|⭐️⭐️⭐️|
+| [[LLM推理优化][Attention优化]📖图解:从Online-Softmax到FlashAttention V1/V2/V3](https://zhuanlan.zhihu.com/p/668888063)|@DefTruth|⭐️⭐️⭐️|
+| [[LLM推理优化][Decoding优化]📖原理&图解FlashDecoding/FlashDecoding++](https://zhuanlan.zhihu.com/p/696075602)|@DefTruth|⭐️⭐️|
+| [[VLM推理优化][LLaVA系列]📖CLIP/LLaVA/LLaVA1.5/VILA笔记: 核心点解析](https://zhuanlan.zhihu.com/p/683137074)|@DefTruth|⭐️⭐️|
+| [[LLM推理优化][Attention优化][万字]📖TensorRT MHA/Myelin vs FlashAttention-2](https://zhuanlan.zhihu.com/p/678873216)|@DefTruth|⭐️⭐️⭐️|
+| [[LLM推理优化][PTX汇编]📖CUDA 12 PTX汇编: PRMT指令详解-通用模式](https://zhuanlan.zhihu.com/p/660630414)|@DefTruth|⭐️|
+| [[LLM推理优化][PTX汇编]📖CUDA 12 PTX汇编: LOP3指令详解](https://zhuanlan.zhihu.com/p/659741469)|@DefTruth|⭐️|
+| [[LLM推理优化][CUDA][3w字]📖高频面试题汇总-大模型手撕CUDA](https://zhuanlan.zhihu.com/p/678903537)|@DefTruth|⭐️⭐️⭐️|
+| [[LLM推理优化][Weight Only]📖WINT8/4-(00): 通俗易懂讲解-快速反量化算法](https://zhuanlan.zhihu.com/p/657072856)|@DefTruth|⭐️⭐️|
+| [[LLM推理优化][Weight Only]📖WINT8/4-(01): PRMT指令详解及FT源码解析](https://zhuanlan.zhihu.com/p/657070837)|@DefTruth|⭐️⭐️|
+| [[LLM推理优化][Weight Only]📖WINT8/4-(02): 快速反量化之INT8转BF16](https://zhuanlan.zhihu.com/p/657073159)|@DefTruth|⭐️⭐️|
+| [[LLM推理优化][Weight Only]📖WINT8/4-(03): LOP3指令详解及INT4转FP16/BF16](https://zhuanlan.zhihu.com/p/657073857)|@DefTruth|⭐️⭐️|
+| [[LLM推理优化][LLM Infra整理]📖100+篇: 大模型推理各方向新发展整理](https://zhuanlan.zhihu.com/p/693680304)|@DefTruth|⭐️⭐️|
+| [[LLM推理优化][LLM Infra整理]📖30+篇: LLM推理论文集-500页PDF](https://zhuanlan.zhihu.com/p/669777159)|@DefTruth|⭐️⭐️|
+| [[LLM推理优化][LLM Infra整理]📖FlashDecoding++: 比FlashDecoding还要快！](https://zhuanlan.zhihu.com/p/665022589)|@DefTruth|⭐️|
+| [[LLM推理优化][LLM Infra整理]📖TensorRT-LLM开源，TensorRT 9.1也来了](https://zhuanlan.zhihu.com/p/662361469)|@DefTruth|⭐️|
+| [[LLM推理优化][LLM Infra整理]📖20+篇: LLM推理论文集-300页PDF](https://zhuanlan.zhihu.com/p/658091768)|@DefTruth|⭐️⭐️|
+| [[LLM推理优化][LLM Infra整理]📖PagedAttention论文新鲜出炉](https://zhuanlan.zhihu.com/p/617015570)|@DefTruth|⭐️|
+| [[推理部署][CV/NLP]📖FastDeploy三行代码搞定150+ CV、NLP模型部署](https://zhuanlan.zhihu.com/p/581326442)|@DefTruth|⭐️|
+| [[推理部署][CV]📖如何在lite.ai.toolkit(3.6k+ stars)中增加您的模型？](https://zhuanlan.zhihu.com/p/523876625)|@DefTruth|⭐️⭐️|
+| [[推理部署][CV]📖美团 YOLOv6 ORT/MNN/TNN/NCNN C++推理部署](https://zhuanlan.zhihu.com/p/533643238)|@DefTruth|⭐️⭐️|
+| [[推理部署][ONNX]📖ONNX推理加速技术文档-杂记](https://zhuanlan.zhihu.com/p/524023964)|@DefTruth|⭐️|
+| [[推理部署][TensorFlow]📖Mac源码编译TensorFlow C++指北](https://zhuanlan.zhihu.com/p/524013615)|@DefTruth|⭐️|
+| [[推理部署][CV]📖1Mb!头部姿态估计: FSANet，一个小而美的模型(C++)](https://zhuanlan.zhihu.com/p/447364201)|@DefTruth|⭐️|
+| [[推理部署][CV]📖opencv+ffmpeg编译打包全解指南](https://zhuanlan.zhihu.com/p/472115312)|@DefTruth|⭐️⭐️|
+| [[推理部署][CV]📖RobustVideoMatting视频抠图静态ONNX模型转换](https://zhuanlan.zhihu.com/p/459088407)|@DefTruth|⭐️|
+| [[推理部署][CV]📖190Kb!SSRNet年龄检测详细解读（含C++工程）](https://zhuanlan.zhihu.com/p/462762797)|@DefTruth|⭐️|
+| [[推理部署][CV]📖MGMatting(CVPR2021)人像抠图C++应用记录](https://zhuanlan.zhihu.com/p/464732042)|@DefTruth|⭐️|
+| [[推理部署][CV]📖超准确人脸检测(带关键点)YOLO5Face C++工程详细记录](https://zhuanlan.zhihu.com/p/461878005)|@DefTruth|⭐️⭐️|
+| [[推理部署][ORT]📖解决: ONNXRuntime(Python) GPU 部署配置记录](https://zhuanlan.zhihu.com/p/457484536)|@DefTruth|⭐️|
+| [[推理部署][CV]📖记录SCRFD(CVPR2021)人脸检测C++工程化(含docker镜像)](https://zhuanlan.zhihu.com/p/455165568)|@DefTruth|⭐️⭐️|
+| [[推理部署][NCNN]📖野路子：记录一个解决onnx转ncnn时op不支持的trick](https://zhuanlan.zhihu.com/p/451446147)|@DefTruth|⭐️|
+| [[推理部署][CV]📖升级版轻量级NanoDet-Plus MNN/TNN/NCNN/ORT C++工程记录](https://zhuanlan.zhihu.com/p/450586647)|@DefTruth|⭐️⭐️|
+| [[推理部署][CV]📖超轻量级NanoDet MNN/TNN/NCNN/ORT C++工程记录](https://zhuanlan.zhihu.com/p/443419387)|@DefTruth|⭐️|
+| [[推理部署][CV]📖详细记录MGMatting之MNN、TNN和ORT C++移植](https://zhuanlan.zhihu.com/p/442949027)|@DefTruth|⭐️⭐️|
+| [[推理部署][CV]📖YOLOX NCNN/MNN/TNN/ONNXRuntime C++工程简记](https://zhuanlan.zhihu.com/p/447364122)|@DefTruth|⭐️|
+| [[推理部署][TNN]📖手动修改YoloX的tnnproto记录-TNN](https://zhuanlan.zhihu.com/p/425668734)|@DefTruth|⭐️|
+| [[推理部署][ORT]📖全网最详细 ONNXRuntime C++/Java/Python 资料！](https://zhuanlan.zhihu.com/p/414317269)|@DefTruth|⭐️|
+| [[推理部署][CV]📖RobustVideoMatting: C++工程化记录-实现篇](https://zhuanlan.zhihu.com/p/413280488)|@DefTruth|⭐️⭐️|
+| [[推理部署][CV]📖RobustVideoMatting: C++工程化记录-应用篇](https://zhuanlan.zhihu.com/p/412491918)|@DefTruth|⭐️⭐️|
+| [[推理部署][ORT]📖ONNXRuntime C++ CMake 工程分析及编译](https://zhuanlan.zhihu.com/p/411887386)|@DefTruth|⭐️⭐️|
+| [[推理部署][ORT]📖如何使用ORT C++ API处理NCHW和NHWC输入？](https://zhuanlan.zhihu.com/p/524230808)|@DefTruth|⭐️|
+| [[推理部署][TNN]📖tnn-convert搭建简记-YOLOP转TNN](https://zhuanlan.zhihu.com/p/431418709)|@DefTruth|⭐️|
+| [[推理部署][CV]📖YOLOP ONNXRuntime C++工程化记录](https://zhuanlan.zhihu.com/p/411651933)|@DefTruth|⭐️⭐️|
+| [[推理部署][NCNN]📖超有用NCNN参考资料整理](https://zhuanlan.zhihu.com/p/449765328)|@DefTruth|⭐️|
+| [[推理部署][MNN]📖超有用MNN参考资料整理](https://zhuanlan.zhihu.com/p/449761992)|@DefTruth|⭐️|
+| [[推理部署][TNN]📖超有用TNN参考资料整理](https://zhuanlan.zhihu.com/p/449769615)|@DefTruth|⭐️|
+| [[推理部署][ONNX]📖超有用ONNX参考资料整理](https://zhuanlan.zhihu.com/p/449773663)|@DefTruth|⭐️|
+| [[推理部署][ONNX]📖超有用ONNX模型结构参考资料整理](https://zhuanlan.zhihu.com/p/449775926)|@DefTruth|⭐️|
+| [[推理部署][OpenCV-DNN]📖超有用OpenCV-DNN参考资料整理](https://zhuanlan.zhihu.com/p/449778377)|@DefTruth|⭐️|
+| [[推理部署][Tensorflow]📖超有用Tensorflow C++工程化知识点](https://zhuanlan.zhihu.com/p/449788027)|@DefTruth|⭐️|
+| [[推理部署][模型转换]📖深度学习模型转换资料整理](https://zhuanlan.zhihu.com/p/449759361)|@DefTruth|⭐️|
+| [[技术随笔][C++][CMake]📖超有用CMake参考资料整理](https://zhuanlan.zhihu.com/p/449779892)|@DefTruth|⭐️⭐️|
+| [[技术随笔][C++][3W字]📖静态链接和静态库实践指北-原理篇](https://zhuanlan.zhihu.com/p/595527528)|@DefTruth|⭐️⭐️⭐️|
+| [[技术随笔][C++]📖Mac下C++内存检查指北(Valgrind VS Asan)](https://zhuanlan.zhihu.com/p/508470880)|@DefTruth|⭐️|
+| [[技术随笔][CV]📖torchlm: 人脸关键点检测库](https://zhuanlan.zhihu.com/p/467211561)|@DefTruth|⭐️⭐️|
+| [[技术随笔][ML]📖《统计学习方法-李航: 笔记-从原理到实现-基于R》](https://zhuanlan.zhihu.com/p/684885595)|@DefTruth|⭐️⭐️|
+| [[技术随笔][Git]📖如何优雅地git clone和git submodule？](https://zhuanlan.zhihu.com/p/639136221)|@DefTruth|⭐️|
+| [[技术随笔][3D]📖人脸重建3D参考资料整理](https://zhuanlan.zhihu.com/p/524034741)|@DefTruth|⭐️|
+| [[技术随笔][3D]📖BlendShapes参考资料整理](https://zhuanlan.zhihu.com/p/524036145)|@DefTruth|⭐️|
+| [[技术随笔][3D]📖从源码安装Pytorch3D详细记录及学习资料](https://zhuanlan.zhihu.com/p/512347464)|@DefTruth|⭐️|
+| [[技术随笔][ML]📖200页:《统计学习方法：李航》笔记 -从原理到实现](https://zhuanlan.zhihu.com/p/461520847)|@DefTruth|⭐️⭐️|
 
 ### 📚 高性能计算与分布式-技术博客推荐 ([©️back👆🏻](#contents))
 

From 706e1c9674aae36f0c604fa0881d7f3c862e1ea1 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Sat, 7 Jun 2025 10:17:52 +0800
Subject: [PATCH 12/57] Update README.md

---
 README.md | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/README.md b/README.md
index 51765337..1a2ddd10 100644
--- a/README.md
+++ b/README.md
@@ -510,7 +510,7 @@ The kernels listed here will guide you through a step-by-step progression, rangi
 |📖 类型-标题|📖 作者| 📖 推荐 |
 |:---|:---|:---|
 | [[Triton编程][基础]📖Triton极简入门: Triton Vector Add](https://zhuanlan.zhihu.com/p/1902778199261291694)|@DefTruth|⭐️⭐️⭐|
-| [[Triton编程][基础]📖Triton Fused Softmax Kernel详解: 从Python源码到PTX分析](https://zhuanlan.zhihu.com/p/1899562146477609112)|@DefTruth|⭐️⭐️⭐|
+| [[Triton编程][基础]📖Triton Fused Softmax Kernel详解: 从Python源码到PTX](https://zhuanlan.zhihu.com/p/1899562146477609112)|@DefTruth|⭐️⭐️⭐|
 | [[Triton编程][基础]📖vLLM Triton Merge Attention States Kernel详解](https://zhuanlan.zhihu.com/p/1904937907703243110)|@DefTruth|⭐️⭐️⭐|
 | [[Triton编程][进阶]📖vLLM Prefix Prefill Triton Kernel图解](https://zhuanlan.zhihu.com/p/695799736)|@DefTruth|⭐️⭐️⭐️|
 | [[张量/序列并行]📖序列并行: BPT、Ring-Attention及Striped-Attention笔记](https://zhuanlan.zhihu.com/p/6456708235)|@DefTruth|⭐️⭐️⭐|
@@ -554,7 +554,7 @@ The kernels listed here will guide you through a step-by-step progression, rangi
 | [[推理部署][ORT]📖解决: ONNXRuntime(Python) GPU 部署配置记录](https://zhuanlan.zhihu.com/p/457484536)|@DefTruth|⭐️|
 | [[推理部署][CV]📖记录SCRFD(CVPR2021)人脸检测C++工程化(含docker镜像)](https://zhuanlan.zhihu.com/p/455165568)|@DefTruth|⭐️⭐️|
 | [[推理部署][NCNN]📖野路子：记录一个解决onnx转ncnn时op不支持的trick](https://zhuanlan.zhihu.com/p/451446147)|@DefTruth|⭐️|
-| [[推理部署][CV]📖升级版轻量级NanoDet-Plus MNN/TNN/NCNN/ORT C++工程记录](https://zhuanlan.zhihu.com/p/450586647)|@DefTruth|⭐️⭐️|
+| [[推理部署][CV]📖升级版NanoDet-Plus MNN/TNN/NCNN/ORT C++工程记录](https://zhuanlan.zhihu.com/p/450586647)|@DefTruth|⭐️⭐️|
 | [[推理部署][CV]📖超轻量级NanoDet MNN/TNN/NCNN/ORT C++工程记录](https://zhuanlan.zhihu.com/p/443419387)|@DefTruth|⭐️|
 | [[推理部署][CV]📖详细记录MGMatting之MNN、TNN和ORT C++移植](https://zhuanlan.zhihu.com/p/442949027)|@DefTruth|⭐️⭐️|
 | [[推理部署][CV]📖YOLOX NCNN/MNN/TNN/ONNXRuntime C++工程简记](https://zhuanlan.zhihu.com/p/447364122)|@DefTruth|⭐️|

From ffa027f36a12a1789d8dd7efab22022a7ff6a475 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Sat, 7 Jun 2025 10:25:01 +0800
Subject: [PATCH 13/57] misc: add hgemv cute index

---
 README.md | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/README.md b/README.md
index 1a2ddd10..227c9234 100644
--- a/README.md
+++ b/README.md
@@ -499,6 +499,9 @@ The kernels listed here will guide you through a step-by-step progression, rangi
 |:---|:---|:---|:---|:---|
 | ✔️ [mat_transpose_cute](./kernels/mat-transpose/mat_transpose_cute.cu)|f32|/|[link](./kernels/mat-transpose/)|⭐️⭐️|
 | ✔️ [flash_attn_cute(naive)](./kernels/flash-attn/cutlass/flash_attn_cute.cu)|f16|f32|[link](./kernels/flash-attn/)|⭐️⭐️⭐️|
+| ✔️ [hgemv_f16_cute_kernel](./kernels/hgemv/hgemv_cute.cu)|f16|f16|[link](./kernels/hgemv/)|⭐️⭐️⭐️|
+| ✔️ [hgemv_f16x8_cute_kernel](./kernels/hgemv/hgemv_cute.cu)|f16|f16|[link](./kernels/hgemv/)|⭐️⭐️⭐️|
+| ✔️ [hgemv_tensor_core_cute_kernel](./kernels/hgemv/hgemv_cute.cu)|f16|f16|[link](./kernels/hgemv/)|⭐️⭐️⭐️|
 | ✔️ [hgemm_mma_stages_swizzle{smem}...cute*](./kernels/hgemm/cutlass/hgemm_mma_stage_tn_cute.cu)|f16|f16|[link](./kernels/hgemm/)|⭐️⭐️⭐️|
 
 ## 📖 100+ 高性能计算与分布式-技术博客

From d527537f864ee95d3bc7a13c5be1b74911a37a21 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Sun, 8 Jun 2025 16:43:11 +0800
Subject: [PATCH 14/57] Update README.md

---
 README.md | 29 +++++++++++++++--------------
 1 file changed, 15 insertions(+), 14 deletions(-)

diff --git a/README.md b/README.md
index 227c9234..87e40d20 100644
--- a/README.md
+++ b/README.md
@@ -34,6 +34,18 @@
   </p>
 </div>
 
+## ©️Citations🎉🎉
+
+```BibTeX
+@misc{LeetCUDA@2025,
+  title={LeetCUDA: A Modern CUDA Learn Notes with PyTorch for Beginners},
+  url={https://github.com/xlite-dev/LeetCUDA.git},
+  note={Open-source software available at https://github.com/xlite-dev/LeetCUDA.git},
+  author={DefTruth and Many Others},
+  year={2025}
+}
+```
+
 
 ## 📖 News 🔥🔥
 <div id="news"></div>
@@ -54,6 +66,7 @@
   <img src='https://github.com/user-attachments/assets/9472e970-c083-4b31-9252-3eeecc761078' height="170px" width="270px">
 </div>
 
+
 ## 📖 Contents
 <div id="contents"></div>
 <!---
@@ -98,7 +111,7 @@
   - [📚 Hard++ ⭐⭐⭐️⭐️⭐️](#cuda-kernel-hard-plus)
   - [📚 Triton ⭐⭐⭐️](#triton-kernel)
   - [📚 CUTLASS ⭐⭐⭐️](#cutlass-kernel)
-- [📖 100+ 高性能计算文章 💡💡](#my-blogs-part-1)
+- [📖 100+ LLM/CUDA Blogs 🔥](#my-blogs-part-1)
 - [📖 How to Contribute 👀👇](#contribute)
 
 
@@ -225,18 +238,6 @@ flash_attn_mma_stages_split_q_tiling_qkv_kernel(half* Q, half* K, half* V, half*
 
 💡NOTE: [📚Split Q + Fully QKV Fine-grained Tiling](#mma-tiling-qkv) has been refactored into 🤖[ffpa-attn](https://github.com/xlite-dev/ffpa-attn).
 
-## ©️Citations🎉🎉
-
-```BibTeX
-@misc{LeetCUDA@2025,
-  title={LeetCUDA: A Modern CUDA Learn Notes with PyTorch for Beginners},
-  url={https://github.com/xlite-dev/LeetCUDA},
-  note={Open-source software available at https://github.com/xlite-dev/LeetCUDA},
-  author={DefTruth etc},
-  year={2025}
-}
-```
-
 ## 📖 200+ CUDA Kernels 🔥🔥 (Easy -> Hard++) ([©️back👆🏻](#contents))
 
 <div id="cuda-kernel"></div>
@@ -481,7 +482,7 @@ The kernels listed here will guide you through a step-by-step progression, rangi
 
 💡NOTE: 🤖[ffpa-attn](https://github.com/xlite-dev/ffpa-attn): 📚FFPA - Yet another Faster Flash Prefill Attention with O(1)🎉SRAM complexity for headdim > 256, **1.8x~3x**🎉faster than SDPA EA: [📈L20 ~1.9x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-l20), [📈 A30 ~1.8x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-a30), [📈3080 ~2.9x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-3080), [📈4090 ~2.1x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-4090).
 
-### 📚 Triton Kernel (OpenAI Triton) ([©️back👆🏻](#cuda-kernel))
+### 📚 Triton Kernel (OpenAI Triton) ⭐️⭐️⭐️ ([©️back👆🏻](#cuda-kernel))
 
 <div id="triton-kernel"></div>
 

From 92feb71561969d05aa44cada5a8e6f6f98018815 Mon Sep 17 00:00:00 2001
From: kitecats <90194592+kitecats@users.noreply.github.com>
Date: Mon, 9 Jun 2025 18:22:21 +0800
Subject: [PATCH 15/57] feat: add a cute & bank-free mat transpose (#334)

* Use 128-bit data loading

* add a bank-free mat transpose vectorize impelment

---------

Co-authored-by: ZKJW037 <ZKJW037@ZKJW-HZ020>
---
 kernels/mat-transpose/README.md             | 824 ++++++++++----------
 kernels/mat-transpose/mat_transpose.cu      |   5 +
 kernels/mat-transpose/mat_transpose.py      |  11 +-
 kernels/mat-transpose/mat_transpose_cute.cu |  94 +++
 4 files changed, 529 insertions(+), 405 deletions(-)

diff --git a/kernels/mat-transpose/README.md b/kernels/mat-transpose/README.md
index f19951aa..8183de38 100755
--- a/kernels/mat-transpose/README.md
+++ b/kernels/mat-transpose/README.md
@@ -40,455 +40,471 @@ python3 mat_transpose.py
 ```bash
 ----------------------------------------------------------------------------------------------------------------------------------
                                                        M=1024, N=1024
-                       out_original: [0.0, 1.0, 1024.0], validate False, time:0.00667048ms
-                    out_f32_col2row: [0.0, 1024.0, 1.0], validate True , time:0.01761174ms
-                    out_f32_row2col: [0.0, 1024.0, 1.0], validate True , time:0.01008821ms
-                out_f32_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.00852585ms
-                out_f32_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.00851846ms
-                  out_f32_diagnonal: [0.0, 1024.0, 1.0], validate True , time:0.00641012ms
-                  out_f32x4_col2row: [0.0, 1024.0, 1.0], validate True , time:0.01742840ms
-                  out_f32x4_row2col: [0.0, 1024.0, 1.0], validate True , time:0.00943899ms
-              out_f32x4_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.00987363ms
-              out_f32x4_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.00499630ms
-       out_f32x4_shared_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.00524426ms
-       out_f32x4_shared_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.00670385ms
-   out_f32x4_shared_bcf_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.00536251ms
-   out_f32x4_shared_bcf_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.00580740ms
- out_mat_transpose_cute_col2row_reg: [0.0, 1024.0, 1.0], validate True , time:0.00747943ms
- out_mat_transpose_cute_row2col_reg: [0.0, 1024.0, 1.0], validate True , time:0.00599384ms
-    out_mat_transpose_cute_col_smem: [0.0, 1024.0, 1.0], validate True , time:0.00989509ms
-    out_mat_transpose_cute_row_smem: [0.0, 1024.0, 1.0], validate True , time:0.00534320ms
-out_mat_transpose_cute_col_smem_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.00869393ms
-out_mat_transpose_cute_row_smem_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.00555682ms
-out_mat_transpose_cute_row_cvectorized: [0.0, 1024.0, 1.0], validate True , time:0.00727987ms
-out_mat_transpose_cute_row_rvectorized: [0.0, 1024.0, 1.0], validate True , time:0.00493646ms
-out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.00592184ms
-out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.00563526ms
-                         out_f32_th: [0.0, 1024.0, 1.0], validate True , time:0.01460934ms
-                out_f32_th_compiled: [0.0, 1024.0, 1.0], validate True , time:0.03173113ms
+                       out_original: [0.0, 1.0, 1024.0], validate False, time:0.01322293ms
+                    out_f32_col2row: [0.0, 1024.0, 1.0], validate True , time:0.03734207ms
+                    out_f32_row2col: [0.0, 1024.0, 1.0], validate True , time:0.02030921ms
+                out_f32_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.01762843ms
+                out_f32_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.01770711ms
+                  out_f32_diagnonal: [0.0, 1024.0, 1.0], validate True , time:0.01362014ms
+                  out_f32x4_col2row: [0.0, 1024.0, 1.0], validate True , time:0.03673577ms
+                  out_f32x4_row2col: [0.0, 1024.0, 1.0], validate True , time:0.01884627ms
+              out_f32x4_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.01892900ms
+              out_f32x4_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.01024675ms
+       out_f32x4_shared_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.01027608ms
+       out_f32x4_shared_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.01155329ms
+   out_f32x4_shared_bcf_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.01036501ms
+   out_f32x4_shared_bcf_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.01045060ms
+ out_mat_transpose_cute_col2row_reg: [0.0, 1024.0, 1.0], validate True , time:0.01586843ms
+ out_mat_transpose_cute_row2col_reg: [0.0, 1024.0, 1.0], validate True , time:0.01302576ms
+    out_mat_transpose_cute_col_smem: [0.0, 1024.0, 1.0], validate True , time:0.01425385ms
+    out_mat_transpose_cute_row_smem: [0.0, 1024.0, 1.0], validate True , time:0.01313591ms
+out_mat_transpose_cute_col_smem_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.01398778ms
+out_mat_transpose_cute_row_smem_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.01285744ms
+out_mat_transpose_cute_row_cvectorized: [0.0, 1024.0, 1.0], validate True , time:0.01234102ms
+out_mat_transpose_cute_row_rvectorized: [0.0, 1024.0, 1.0], validate True , time:0.01036048ms
+out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.01147723ms
+out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.01037669ms
+out_mat_transpose_cute_row_rvectorized_swizzled_optimized: [0.0, 1024.0, 1.0], validate True , time:0.01087165ms
+                         out_f32_th: [0.0, 1024.0, 1.0], validate True , time:0.03082085ms
+                out_f32_th_compiled: [0.0, 1024.0, 1.0], validate True , time:0.07333684ms
 ----------------------------------------------------------------------------------------------------------------------------------
 ----------------------------------------------------------------------------------------------------------------------------------
                                                        M=1024, N=2048
-                       out_original: [0.0, 1.0, 2048.0], validate False, time:0.00685191ms
-                    out_f32_col2row: [0.0, 2048.0, 1.0], validate True , time:0.03323030ms
-                    out_f32_row2col: [0.0, 2048.0, 1.0], validate True , time:0.01826382ms
-                out_f32_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.01468587ms
-                out_f32_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.01467657ms
-                  out_f32x4_col2row: [0.0, 2048.0, 1.0], validate True , time:0.03329539ms
-                  out_f32x4_row2col: [0.0, 2048.0, 1.0], validate True , time:0.01831269ms
-              out_f32x4_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.01730418ms
-              out_f32x4_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.00699377ms
-       out_f32x4_shared_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.00813270ms
-       out_f32x4_shared_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.01007080ms
-   out_f32x4_shared_bcf_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.00836945ms
-   out_f32x4_shared_bcf_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.00849128ms
- out_mat_transpose_cute_col2row_reg: [0.0, 2048.0, 1.0], validate True , time:0.01324391ms
- out_mat_transpose_cute_row2col_reg: [0.0, 2048.0, 1.0], validate True , time:0.00926590ms
-    out_mat_transpose_cute_col_smem: [0.0, 2048.0, 1.0], validate True , time:0.01706100ms
-    out_mat_transpose_cute_row_smem: [0.0, 2048.0, 1.0], validate True , time:0.00816727ms
-out_mat_transpose_cute_col_smem_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.01489687ms
-out_mat_transpose_cute_row_smem_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.00876975ms
-out_mat_transpose_cute_row_cvectorized: [0.0, 2048.0, 1.0], validate True , time:0.01191163ms
-out_mat_transpose_cute_row_rvectorized: [0.0, 2048.0, 1.0], validate True , time:0.00748754ms
-out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.00862432ms
-out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.00858784ms
-                         out_f32_th: [0.0, 2048.0, 1.0], validate True , time:0.02438354ms
-                out_f32_th_compiled: [0.0, 2048.0, 1.0], validate True , time:0.04902887ms
+                       out_original: [0.0, 1.0, 2048.0], validate False, time:0.02374721ms
+                    out_f32_col2row: [0.0, 2048.0, 1.0], validate True , time:0.07149410ms
+                    out_f32_row2col: [0.0, 2048.0, 1.0], validate True , time:0.03805780ms
+                out_f32_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.03402877ms
+                out_f32_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.03385830ms
+                  out_f32x4_col2row: [0.0, 2048.0, 1.0], validate True , time:0.06924939ms
+                  out_f32x4_row2col: [0.0, 2048.0, 1.0], validate True , time:0.03661728ms
+              out_f32x4_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.03592300ms
+              out_f32x4_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.02259111ms
+       out_f32x4_shared_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.02224588ms
+       out_f32x4_shared_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.02293062ms
+   out_f32x4_shared_bcf_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.02227783ms
+   out_f32x4_shared_bcf_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.02276349ms
+ out_mat_transpose_cute_col2row_reg: [0.0, 2048.0, 1.0], validate True , time:0.03074288ms
+ out_mat_transpose_cute_row2col_reg: [0.0, 2048.0, 1.0], validate True , time:0.02413964ms
+    out_mat_transpose_cute_col_smem: [0.0, 2048.0, 1.0], validate True , time:0.02684236ms
+    out_mat_transpose_cute_row_smem: [0.0, 2048.0, 1.0], validate True , time:0.02411056ms
+out_mat_transpose_cute_col_smem_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.02621555ms
+out_mat_transpose_cute_row_smem_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.02390814ms
+out_mat_transpose_cute_row_cvectorized: [0.0, 2048.0, 1.0], validate True , time:0.02338839ms
+out_mat_transpose_cute_row_rvectorized: [0.0, 2048.0, 1.0], validate True , time:0.02213287ms
+out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.02302265ms
+out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.02224278ms
+out_mat_transpose_cute_row_rvectorized_swizzled_optimized: [0.0, 2048.0, 1.0], validate True , time:0.02258778ms
+                         out_f32_th: [0.0, 2048.0, 1.0], validate True , time:0.06115103ms
+                out_f32_th_compiled: [0.0, 2048.0, 1.0], validate True , time:0.05778503ms
 ----------------------------------------------------------------------------------------------------------------------------------
 ----------------------------------------------------------------------------------------------------------------------------------
                                                        M=1024, N=4096
-                       out_original: [0.0, 1.0, 4096.0], validate False, time:0.01132798ms
-                    out_f32_col2row: [0.0, 4096.0, 1.0], validate True , time:0.06182218ms
-                    out_f32_row2col: [0.0, 4096.0, 1.0], validate True , time:0.03392482ms
-                out_f32_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.02772188ms
-                out_f32_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.02774906ms
-                  out_f32x4_col2row: [0.0, 4096.0, 1.0], validate True , time:0.06245613ms
-                  out_f32x4_row2col: [0.0, 4096.0, 1.0], validate True , time:0.03352976ms
-              out_f32x4_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.03137302ms
-              out_f32x4_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.01143122ms
-       out_f32x4_shared_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.01373696ms
-       out_f32x4_shared_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.01705766ms
-   out_f32x4_shared_bcf_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.01421714ms
-   out_f32x4_shared_bcf_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.01388431ms
- out_mat_transpose_cute_col2row_reg: [0.0, 4096.0, 1.0], validate True , time:0.02740574ms
- out_mat_transpose_cute_row2col_reg: [0.0, 4096.0, 1.0], validate True , time:0.01666737ms
-    out_mat_transpose_cute_col_smem: [0.0, 4096.0, 1.0], validate True , time:0.03182530ms
-    out_mat_transpose_cute_row_smem: [0.0, 4096.0, 1.0], validate True , time:0.01621008ms
-out_mat_transpose_cute_col_smem_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.02744722ms
-out_mat_transpose_cute_row_smem_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.01582789ms
-out_mat_transpose_cute_row_cvectorized: [0.0, 4096.0, 1.0], validate True , time:0.02123284ms
-out_mat_transpose_cute_row_rvectorized: [0.0, 4096.0, 1.0], validate True , time:0.01262116ms
-out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.01439214ms
-out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.01478362ms
-                         out_f32_th: [0.0, 4096.0, 1.0], validate True , time:0.04604459ms
-                out_f32_th_compiled: [0.0, 4096.0, 1.0], validate True , time:0.04802442ms
+                       out_original: [0.0, 1.0, 4096.0], validate False, time:0.04392624ms
+                    out_f32_col2row: [0.0, 4096.0, 1.0], validate True , time:0.12588978ms
+                    out_f32_row2col: [0.0, 4096.0, 1.0], validate True , time:0.07158637ms
+                out_f32_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.06782699ms
+                out_f32_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.06651688ms
+                  out_f32x4_col2row: [0.0, 4096.0, 1.0], validate True , time:0.13267922ms
+                  out_f32x4_row2col: [0.0, 4096.0, 1.0], validate True , time:0.07224798ms
+              out_f32x4_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.06865716ms
+              out_f32x4_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.04291391ms
+       out_f32x4_shared_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.04286146ms
+       out_f32x4_shared_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.04337263ms
+   out_f32x4_shared_bcf_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.04302001ms
+   out_f32x4_shared_bcf_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.04323077ms
+ out_mat_transpose_cute_col2row_reg: [0.0, 4096.0, 1.0], validate True , time:0.06857824ms
+ out_mat_transpose_cute_row2col_reg: [0.0, 4096.0, 1.0], validate True , time:0.04765439ms
+    out_mat_transpose_cute_col_smem: [0.0, 4096.0, 1.0], validate True , time:0.05450273ms
+    out_mat_transpose_cute_row_smem: [0.0, 4096.0, 1.0], validate True , time:0.04703522ms
+out_mat_transpose_cute_col_smem_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.05180621ms
+out_mat_transpose_cute_row_smem_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.04637504ms
+out_mat_transpose_cute_row_cvectorized: [0.0, 4096.0, 1.0], validate True , time:0.04552484ms
+out_mat_transpose_cute_row_rvectorized: [0.0, 4096.0, 1.0], validate True , time:0.04266071ms
+out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.04476309ms
+out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.04267573ms
+out_mat_transpose_cute_row_rvectorized_swizzled_optimized: [0.0, 4096.0, 1.0], validate True , time:0.04308534ms
+                         out_f32_th: [0.0, 4096.0, 1.0], validate True , time:0.11653090ms
+                out_f32_th_compiled: [0.0, 4096.0, 1.0], validate True , time:0.11728549ms
 ----------------------------------------------------------------------------------------------------------------------------------
 ----------------------------------------------------------------------------------------------------------------------------------
                                                        M=1024, N=8192
-                       out_original: [0.0, 1.0, 8192.0], validate False, time:0.02622271ms
-                    out_f32_col2row: [0.0, 8192.0, 1.0], validate True , time:0.13227582ms
-                    out_f32_row2col: [0.0, 8192.0, 1.0], validate True , time:0.06796670ms
-                out_f32_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.06429243ms
-                out_f32_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.06427002ms
-                  out_f32x4_col2row: [0.0, 8192.0, 1.0], validate True , time:0.13971567ms
-                  out_f32x4_row2col: [0.0, 8192.0, 1.0], validate True , time:0.06687403ms
-              out_f32x4_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.08264017ms
-              out_f32x4_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.02314472ms
-       out_f32x4_shared_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.02923679ms
-       out_f32x4_shared_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.03703094ms
-   out_f32x4_shared_bcf_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.02966428ms
-   out_f32x4_shared_bcf_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.03264546ms
- out_mat_transpose_cute_col2row_reg: [0.0, 8192.0, 1.0], validate True , time:0.05948019ms
- out_mat_transpose_cute_row2col_reg: [0.0, 8192.0, 1.0], validate True , time:0.04069352ms
-    out_mat_transpose_cute_col_smem: [0.0, 8192.0, 1.0], validate True , time:0.06174302ms
-    out_mat_transpose_cute_row_smem: [0.0, 8192.0, 1.0], validate True , time:0.04106617ms
-out_mat_transpose_cute_col_smem_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.05320787ms
-out_mat_transpose_cute_row_smem_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.03881240ms
-out_mat_transpose_cute_row_cvectorized: [0.0, 8192.0, 1.0], validate True , time:0.04451799ms
-out_mat_transpose_cute_row_rvectorized: [0.0, 8192.0, 1.0], validate True , time:0.02818966ms
-out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.03275704ms
-out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.03005648ms
-                         out_f32_th: [0.0, 8192.0, 1.0], validate True , time:0.09130526ms
-                out_f32_th_compiled: [0.0, 8192.0, 1.0], validate True , time:0.04808426ms
+                       out_original: [0.0, 1.0, 8192.0], validate False, time:0.08478189ms
+                    out_f32_col2row: [0.0, 8192.0, 1.0], validate True , time:0.25058746ms
+                    out_f32_row2col: [0.0, 8192.0, 1.0], validate True , time:0.14558482ms
+                out_f32_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.13590932ms
+                out_f32_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.13259697ms
+                  out_f32x4_col2row: [0.0, 8192.0, 1.0], validate True , time:0.26455188ms
+                  out_f32x4_row2col: [0.0, 8192.0, 1.0], validate True , time:0.16565704ms
+              out_f32x4_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.13458133ms
+              out_f32x4_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.08459783ms
+       out_f32x4_shared_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.08723879ms
+       out_f32x4_shared_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.08400655ms
+   out_f32x4_shared_bcf_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.08733034ms
+   out_f32x4_shared_bcf_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.08398318ms
+ out_mat_transpose_cute_col2row_reg: [0.0, 8192.0, 1.0], validate True , time:0.13130641ms
+ out_mat_transpose_cute_row2col_reg: [0.0, 8192.0, 1.0], validate True , time:0.09233975ms
+    out_mat_transpose_cute_col_smem: [0.0, 8192.0, 1.0], validate True , time:0.10645866ms
+    out_mat_transpose_cute_row_smem: [0.0, 8192.0, 1.0], validate True , time:0.09159970ms
+out_mat_transpose_cute_col_smem_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.10360169ms
+out_mat_transpose_cute_row_smem_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.09012818ms
+out_mat_transpose_cute_row_cvectorized: [0.0, 8192.0, 1.0], validate True , time:0.09122086ms
+out_mat_transpose_cute_row_rvectorized: [0.0, 8192.0, 1.0], validate True , time:0.08314061ms
+out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.08987474ms
+out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.08318615ms
+out_mat_transpose_cute_row_rvectorized_swizzled_optimized: [0.0, 8192.0, 1.0], validate True , time:0.08382916ms
+                         out_f32_th: [0.0, 8192.0, 1.0], validate True , time:0.23323345ms
+                out_f32_th_compiled: [0.0, 8192.0, 1.0], validate True , time:0.23502731ms
 ----------------------------------------------------------------------------------------------------------------------------------
 ----------------------------------------------------------------------------------------------------------------------------------
                                                        M=2048, N=1024
-                       out_original: [0.0, 1.0, 1024.0], validate False, time:0.00638318ms
-                    out_f32_col2row: [0.0, 1024.0, 1.0], validate True , time:0.03563547ms
-                    out_f32_row2col: [0.0, 1024.0, 1.0], validate True , time:0.01650119ms
-                out_f32_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.01620793ms
-                out_f32_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.01620555ms
-                  out_f32x4_col2row: [0.0, 1024.0, 1.0], validate True , time:0.03585052ms
-                  out_f32x4_row2col: [0.0, 1024.0, 1.0], validate True , time:0.01708913ms
-              out_f32x4_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.01859522ms
-              out_f32x4_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.00708890ms
-       out_f32x4_shared_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.00807047ms
-       out_f32x4_shared_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.01017618ms
-   out_f32x4_shared_bcf_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.00838590ms
-   out_f32x4_shared_bcf_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.00867391ms
- out_mat_transpose_cute_col2row_reg: [0.0, 1024.0, 1.0], validate True , time:0.01354027ms
- out_mat_transpose_cute_row2col_reg: [0.0, 1024.0, 1.0], validate True , time:0.00928712ms
-    out_mat_transpose_cute_col_smem: [0.0, 1024.0, 1.0], validate True , time:0.01707125ms
-    out_mat_transpose_cute_row_smem: [0.0, 1024.0, 1.0], validate True , time:0.00816989ms
-out_mat_transpose_cute_col_smem_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.01489687ms
-out_mat_transpose_cute_row_smem_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.00876665ms
-out_mat_transpose_cute_row_cvectorized: [0.0, 1024.0, 1.0], validate True , time:0.01165247ms
-out_mat_transpose_cute_row_rvectorized: [0.0, 1024.0, 1.0], validate True , time:0.00741124ms
-out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.00810361ms
-out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.00862789ms
-                         out_f32_th: [0.0, 1024.0, 1.0], validate True , time:0.02319264ms
-                out_f32_th_compiled: [0.0, 1024.0, 1.0], validate True , time:0.05861235ms
+                       out_original: [0.0, 1.0, 1024.0], validate False, time:0.02391768ms
+                    out_f32_col2row: [0.0, 1024.0, 1.0], validate True , time:0.06568789ms
+                    out_f32_row2col: [0.0, 1024.0, 1.0], validate True , time:0.03711987ms
+                out_f32_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.03264403ms
+                out_f32_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.03245735ms
+                  out_f32x4_col2row: [0.0, 1024.0, 1.0], validate True , time:0.06936216ms
+                  out_f32x4_row2col: [0.0, 1024.0, 1.0], validate True , time:0.03759336ms
+              out_f32x4_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.03599405ms
+              out_f32x4_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.02248406ms
+       out_f32x4_shared_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.02231026ms
+       out_f32x4_shared_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.02310872ms
+   out_f32x4_shared_bcf_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.02236223ms
+   out_f32x4_shared_bcf_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.02289176ms
+ out_mat_transpose_cute_col2row_reg: [0.0, 1024.0, 1.0], validate True , time:0.03009558ms
+ out_mat_transpose_cute_row2col_reg: [0.0, 1024.0, 1.0], validate True , time:0.02510953ms
+    out_mat_transpose_cute_col_smem: [0.0, 1024.0, 1.0], validate True , time:0.02802110ms
+    out_mat_transpose_cute_row_smem: [0.0, 1024.0, 1.0], validate True , time:0.02494168ms
+out_mat_transpose_cute_col_smem_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.02729082ms
+out_mat_transpose_cute_row_smem_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.02456188ms
+out_mat_transpose_cute_row_cvectorized: [0.0, 1024.0, 1.0], validate True , time:0.02376342ms
+out_mat_transpose_cute_row_rvectorized: [0.0, 1024.0, 1.0], validate True , time:0.02220488ms
+out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.02344656ms
+out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.02236819ms
+out_mat_transpose_cute_row_rvectorized_swizzled_optimized: [0.0, 1024.0, 1.0], validate True , time:0.02258658ms
+                         out_f32_th: [0.0, 1024.0, 1.0], validate True , time:0.06107378ms
+                out_f32_th_compiled: [0.0, 1024.0, 1.0], validate True , time:0.06182599ms
 ----------------------------------------------------------------------------------------------------------------------------------
 ----------------------------------------------------------------------------------------------------------------------------------
                                                        M=2048, N=2048
-                       out_original: [0.0, 1.0, 2048.0], validate False, time:0.01117110ms
-                    out_f32_col2row: [0.0, 2048.0, 1.0], validate True , time:0.06950235ms
-                    out_f32_row2col: [0.0, 2048.0, 1.0], validate True , time:0.03485703ms
-                out_f32_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.03005457ms
-                out_f32_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.03008485ms
-                  out_f32_diagnonal: [0.0, 2048.0, 1.0], validate True , time:0.01830244ms
-                  out_f32x4_col2row: [0.0, 2048.0, 1.0], validate True , time:0.07024670ms
-                  out_f32x4_row2col: [0.0, 2048.0, 1.0], validate True , time:0.03537869ms
-              out_f32x4_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.03516817ms
-              out_f32x4_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.01145554ms
-       out_f32x4_shared_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.01382899ms
-       out_f32x4_shared_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.01729012ms
-   out_f32x4_shared_bcf_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.01426935ms
-   out_f32x4_shared_bcf_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.01422572ms
- out_mat_transpose_cute_col2row_reg: [0.0, 2048.0, 1.0], validate True , time:0.02737832ms
- out_mat_transpose_cute_row2col_reg: [0.0, 2048.0, 1.0], validate True , time:0.01676321ms
-    out_mat_transpose_cute_col_smem: [0.0, 2048.0, 1.0], validate True , time:0.03183579ms
-    out_mat_transpose_cute_row_smem: [0.0, 2048.0, 1.0], validate True , time:0.01636028ms
-out_mat_transpose_cute_col_smem_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.02742934ms
-out_mat_transpose_cute_row_smem_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.01592755ms
-out_mat_transpose_cute_row_cvectorized: [0.0, 2048.0, 1.0], validate True , time:0.02093101ms
-out_mat_transpose_cute_row_rvectorized: [0.0, 2048.0, 1.0], validate True , time:0.01259851ms
-out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.01370025ms
-out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.01482153ms
-                         out_f32_th: [0.0, 2048.0, 1.0], validate True , time:0.04773760ms
-                out_f32_th_compiled: [0.0, 2048.0, 1.0], validate True , time:0.05919909ms
+                       out_original: [0.0, 1.0, 2048.0], validate False, time:0.04422760ms
+                    out_f32_col2row: [0.0, 2048.0, 1.0], validate True , time:0.12561965ms
+                    out_f32_row2col: [0.0, 2048.0, 1.0], validate True , time:0.07676911ms
+                out_f32_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.06741023ms
+                out_f32_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.06514311ms
+                  out_f32_diagnonal: [0.0, 2048.0, 1.0], validate True , time:0.05187368ms
+                  out_f32x4_col2row: [0.0, 2048.0, 1.0], validate True , time:0.13291311ms
+                  out_f32x4_row2col: [0.0, 2048.0, 1.0], validate True , time:0.07298565ms
+              out_f32x4_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.06908822ms
+              out_f32x4_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.04279065ms
+       out_f32x4_shared_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.04279661ms
+       out_f32x4_shared_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.04364467ms
+   out_f32x4_shared_bcf_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.04293537ms
+   out_f32x4_shared_bcf_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.04339409ms
+ out_mat_transpose_cute_col2row_reg: [0.0, 2048.0, 1.0], validate True , time:0.06557107ms
+ out_mat_transpose_cute_row2col_reg: [0.0, 2048.0, 1.0], validate True , time:0.05029869ms
+    out_mat_transpose_cute_col_smem: [0.0, 2048.0, 1.0], validate True , time:0.05644822ms
+    out_mat_transpose_cute_row_smem: [0.0, 2048.0, 1.0], validate True , time:0.04912710ms
+out_mat_transpose_cute_col_smem_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.05456066ms
+out_mat_transpose_cute_row_smem_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.04846501ms
+out_mat_transpose_cute_row_cvectorized: [0.0, 2048.0, 1.0], validate True , time:0.04663944ms
+out_mat_transpose_cute_row_rvectorized: [0.0, 2048.0, 1.0], validate True , time:0.04294229ms
+out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.04629183ms
+out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.04302597ms
+out_mat_transpose_cute_row_rvectorized_swizzled_optimized: [0.0, 2048.0, 1.0], validate True , time:0.04354882ms
+                         out_f32_th: [0.0, 2048.0, 1.0], validate True , time:0.11834693ms
+                out_f32_th_compiled: [0.0, 2048.0, 1.0], validate True , time:0.11902070ms
 ----------------------------------------------------------------------------------------------------------------------------------
 ----------------------------------------------------------------------------------------------------------------------------------
                                                        M=2048, N=4096
-                       out_original: [0.0, 1.0, 4096.0], validate False, time:0.02612376ms
-                    out_f32_col2row: [0.0, 4096.0, 1.0], validate True , time:0.13792729ms
-                    out_f32_row2col: [0.0, 4096.0, 1.0], validate True , time:0.06782484ms
-                out_f32_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.06800270ms
-                out_f32_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.06801772ms
-                  out_f32x4_col2row: [0.0, 4096.0, 1.0], validate True , time:0.14473867ms
-                  out_f32x4_row2col: [0.0, 4096.0, 1.0], validate True , time:0.06878757ms
-              out_f32x4_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.08687305ms
-              out_f32x4_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.02311611ms
-       out_f32x4_shared_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.02917051ms
-       out_f32x4_shared_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.03698635ms
-   out_f32x4_shared_bcf_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.02948761ms
-   out_f32x4_shared_bcf_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.03269577ms
- out_mat_transpose_cute_col2row_reg: [0.0, 4096.0, 1.0], validate True , time:0.05988336ms
- out_mat_transpose_cute_row2col_reg: [0.0, 4096.0, 1.0], validate True , time:0.04102373ms
-    out_mat_transpose_cute_col_smem: [0.0, 4096.0, 1.0], validate True , time:0.06177592ms
-    out_mat_transpose_cute_row_smem: [0.0, 4096.0, 1.0], validate True , time:0.04124475ms
-out_mat_transpose_cute_col_smem_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.05327010ms
-out_mat_transpose_cute_row_smem_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.03905988ms
-out_mat_transpose_cute_row_cvectorized: [0.0, 4096.0, 1.0], validate True , time:0.04370165ms
-out_mat_transpose_cute_row_rvectorized: [0.0, 4096.0, 1.0], validate True , time:0.02813411ms
-out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.03098154ms
-out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.03004074ms
-                         out_f32_th: [0.0, 4096.0, 1.0], validate True , time:0.09264803ms
-                out_f32_th_compiled: [0.0, 4096.0, 1.0], validate True , time:0.05882740ms
+                       out_original: [0.0, 1.0, 4096.0], validate False, time:0.08467579ms
+                    out_f32_col2row: [0.0, 4096.0, 1.0], validate True , time:0.24970269ms
+                    out_f32_row2col: [0.0, 4096.0, 1.0], validate True , time:0.15055704ms
+                out_f32_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.13335991ms
+                out_f32_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.13127899ms
+                  out_f32x4_col2row: [0.0, 4096.0, 1.0], validate True , time:0.26483536ms
+                  out_f32x4_row2col: [0.0, 4096.0, 1.0], validate True , time:0.15927362ms
+              out_f32x4_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.13693118ms
+              out_f32x4_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.08926201ms
+       out_f32x4_shared_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.09092736ms
+       out_f32x4_shared_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.08591580ms
+   out_f32x4_shared_bcf_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.09117699ms
+   out_f32x4_shared_bcf_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.08632779ms
+ out_mat_transpose_cute_col2row_reg: [0.0, 4096.0, 1.0], validate True , time:0.12909770ms
+ out_mat_transpose_cute_row2col_reg: [0.0, 4096.0, 1.0], validate True , time:0.09749031ms
+    out_mat_transpose_cute_col_smem: [0.0, 4096.0, 1.0], validate True , time:0.11034966ms
+    out_mat_transpose_cute_row_smem: [0.0, 4096.0, 1.0], validate True , time:0.09659719ms
+out_mat_transpose_cute_col_smem_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.10805702ms
+out_mat_transpose_cute_row_smem_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.09507823ms
+out_mat_transpose_cute_row_cvectorized: [0.0, 4096.0, 1.0], validate True , time:0.09092307ms
+out_mat_transpose_cute_row_rvectorized: [0.0, 4096.0, 1.0], validate True , time:0.08320594ms
+out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.09008026ms
+out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.08336163ms
+out_mat_transpose_cute_row_rvectorized_swizzled_optimized: [0.0, 4096.0, 1.0], validate True , time:0.08401752ms
+                         out_f32_th: [0.0, 4096.0, 1.0], validate True , time:0.23449445ms
+                out_f32_th_compiled: [0.0, 4096.0, 1.0], validate True , time:0.23755503ms
 ----------------------------------------------------------------------------------------------------------------------------------
 ----------------------------------------------------------------------------------------------------------------------------------
                                                        M=2048, N=8192
-                       out_original: [0.0, 1.0, 8192.0], validate False, time:0.04788351ms
-                    out_f32_col2row: [0.0, 8192.0, 1.0], validate True , time:0.26702380ms
-                    out_f32_row2col: [0.0, 8192.0, 1.0], validate True , time:0.12641311ms
-                out_f32_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.12766623ms
-                out_f32_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.12762928ms
-                  out_f32x4_col2row: [0.0, 8192.0, 1.0], validate True , time:0.27893376ms
-                  out_f32x4_row2col: [0.0, 8192.0, 1.0], validate True , time:0.12813592ms
-              out_f32x4_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.16845822ms
-              out_f32x4_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.04240894ms
-       out_f32x4_shared_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.05404329ms
-       out_f32x4_shared_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.06824255ms
-   out_f32x4_shared_bcf_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.05481243ms
-   out_f32x4_shared_bcf_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.06005955ms
- out_mat_transpose_cute_col2row_reg: [0.0, 8192.0, 1.0], validate True , time:0.11488175ms
- out_mat_transpose_cute_row2col_reg: [0.0, 8192.0, 1.0], validate True , time:0.07866359ms
-    out_mat_transpose_cute_col_smem: [0.0, 8192.0, 1.0], validate True , time:0.11966443ms
-    out_mat_transpose_cute_row_smem: [0.0, 8192.0, 1.0], validate True , time:0.07921481ms
-out_mat_transpose_cute_col_smem_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.10269117ms
-out_mat_transpose_cute_row_smem_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.07479906ms
-out_mat_transpose_cute_row_cvectorized: [0.0, 8192.0, 1.0], validate True , time:0.08218861ms
-out_mat_transpose_cute_row_rvectorized: [0.0, 8192.0, 1.0], validate True , time:0.05100894ms
-out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.05682635ms
-out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.05597353ms
-                         out_f32_th: [0.0, 8192.0, 1.0], validate True , time:0.17199111ms
-                out_f32_th_compiled: [0.0, 8192.0, 1.0], validate True , time:0.05846906ms
+                       out_original: [0.0, 1.0, 8192.0], validate False, time:0.16631269ms
+                    out_f32_col2row: [0.0, 8192.0, 1.0], validate True , time:0.50284362ms
+                    out_f32_row2col: [0.0, 8192.0, 1.0], validate True , time:0.32336164ms
+                out_f32_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.27494860ms
+                out_f32_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.26218605ms
+                  out_f32x4_col2row: [0.0, 8192.0, 1.0], validate True , time:0.51392102ms
+                  out_f32x4_row2col: [0.0, 8192.0, 1.0], validate True , time:0.37878942ms
+              out_f32x4_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.26178098ms
+              out_f32x4_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.16887903ms
+       out_f32x4_shared_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.17632246ms
+       out_f32x4_shared_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.16928768ms
+   out_f32x4_shared_bcf_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.17651772ms
+   out_f32x4_shared_bcf_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.16609192ms
+ out_mat_transpose_cute_col2row_reg: [0.0, 8192.0, 1.0], validate True , time:0.25982738ms
+ out_mat_transpose_cute_row2col_reg: [0.0, 8192.0, 1.0], validate True , time:0.19544315ms
+    out_mat_transpose_cute_col_smem: [0.0, 8192.0, 1.0], validate True , time:0.21964002ms
+    out_mat_transpose_cute_row_smem: [0.0, 8192.0, 1.0], validate True , time:0.19438195ms
+out_mat_transpose_cute_col_smem_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.21328235ms
+out_mat_transpose_cute_row_smem_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.19120145ms
+out_mat_transpose_cute_row_cvectorized: [0.0, 8192.0, 1.0], validate True , time:0.19285893ms
+out_mat_transpose_cute_row_rvectorized: [0.0, 8192.0, 1.0], validate True , time:0.16419482ms
+out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.19116139ms
+out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.16445112ms
+out_mat_transpose_cute_row_rvectorized_swizzled_optimized: [0.0, 8192.0, 1.0], validate True , time:0.16537309ms
+                         out_f32_th: [0.0, 8192.0, 1.0], validate True , time:0.46839619ms
+                out_f32_th_compiled: [0.0, 8192.0, 1.0], validate True , time:0.46951747ms
 ----------------------------------------------------------------------------------------------------------------------------------
 ----------------------------------------------------------------------------------------------------------------------------------
                                                        M=4096, N=1024
-                       out_original: [0.0, 1.0, 1024.0], validate False, time:0.01135612ms
-                    out_f32_col2row: [0.0, 1024.0, 1.0], validate True , time:0.06473422ms
-                    out_f32_row2col: [0.0, 1024.0, 1.0], validate True , time:0.03022242ms
-                out_f32_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.02885413ms
-                out_f32_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.02884769ms
-                  out_f32x4_col2row: [0.0, 1024.0, 1.0], validate True , time:0.06478477ms
-                  out_f32x4_row2col: [0.0, 1024.0, 1.0], validate True , time:0.03187799ms
-              out_f32x4_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.03276491ms
-              out_f32x4_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.01148319ms
-       out_f32x4_shared_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.01368475ms
-       out_f32x4_shared_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.01800251ms
-   out_f32x4_shared_bcf_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.01427579ms
-   out_f32x4_shared_bcf_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.01484346ms
- out_mat_transpose_cute_col2row_reg: [0.0, 1024.0, 1.0], validate True , time:0.02777314ms
- out_mat_transpose_cute_row2col_reg: [0.0, 1024.0, 1.0], validate True , time:0.01661682ms
-    out_mat_transpose_cute_col_smem: [0.0, 1024.0, 1.0], validate True , time:0.03186488ms
-    out_mat_transpose_cute_row_smem: [0.0, 1024.0, 1.0], validate True , time:0.01607609ms
-out_mat_transpose_cute_col_smem_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.02742004ms
-out_mat_transpose_cute_row_smem_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.01581359ms
-out_mat_transpose_cute_row_cvectorized: [0.0, 1024.0, 1.0], validate True , time:0.02088809ms
-out_mat_transpose_cute_row_rvectorized: [0.0, 1024.0, 1.0], validate True , time:0.01251507ms
-out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.01367640ms
-out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.01482892ms
-                         out_f32_th: [0.0, 1024.0, 1.0], validate True , time:0.04351592ms
-                out_f32_th_compiled: [0.0, 1024.0, 1.0], validate True , time:0.06099153ms
+                       out_original: [0.0, 1.0, 1024.0], validate False, time:0.04421234ms
+                    out_f32_col2row: [0.0, 1024.0, 1.0], validate True , time:0.12696767ms
+                    out_f32_row2col: [0.0, 1024.0, 1.0], validate True , time:0.07942963ms
+                out_f32_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.06330681ms
+                out_f32_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.06264949ms
+                  out_f32x4_col2row: [0.0, 1024.0, 1.0], validate True , time:0.13269401ms
+                  out_f32x4_row2col: [0.0, 1024.0, 1.0], validate True , time:0.07409143ms
+              out_f32x4_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.07032919ms
+              out_f32x4_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.04281425ms
+       out_f32x4_shared_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.04285693ms
+       out_f32x4_shared_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.04353714ms
+   out_f32x4_shared_bcf_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.04288602ms
+   out_f32x4_shared_bcf_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.04316950ms
+ out_mat_transpose_cute_col2row_reg: [0.0, 1024.0, 1.0], validate True , time:0.06419706ms
+ out_mat_transpose_cute_row2col_reg: [0.0, 1024.0, 1.0], validate True , time:0.05924821ms
+    out_mat_transpose_cute_col_smem: [0.0, 1024.0, 1.0], validate True , time:0.06122255ms
+    out_mat_transpose_cute_row_smem: [0.0, 1024.0, 1.0], validate True , time:0.05818987ms
+out_mat_transpose_cute_col_smem_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.06006002ms
+out_mat_transpose_cute_row_smem_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.05800104ms
+out_mat_transpose_cute_row_cvectorized: [0.0, 1024.0, 1.0], validate True , time:0.04861760ms
+out_mat_transpose_cute_row_rvectorized: [0.0, 1024.0, 1.0], validate True , time:0.04257965ms
+out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.04820013ms
+out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.04268837ms
+out_mat_transpose_cute_row_rvectorized_swizzled_optimized: [0.0, 1024.0, 1.0], validate True , time:0.04290032ms
+                         out_f32_th: [0.0, 1024.0, 1.0], validate True , time:0.12000775ms
+                out_f32_th_compiled: [0.0, 1024.0, 1.0], validate True , time:0.12048841ms
 ----------------------------------------------------------------------------------------------------------------------------------
 ----------------------------------------------------------------------------------------------------------------------------------
                                                        M=4096, N=2048
-                       out_original: [0.0, 1.0, 2048.0], validate False, time:0.02611160ms
-                    out_f32_col2row: [0.0, 2048.0, 1.0], validate True , time:0.14095616ms
-                    out_f32_row2col: [0.0, 2048.0, 1.0], validate True , time:0.06800532ms
-                out_f32_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.06891775ms
-                out_f32_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.06890607ms
-                  out_f32x4_col2row: [0.0, 2048.0, 1.0], validate True , time:0.14629245ms
-                  out_f32x4_row2col: [0.0, 2048.0, 1.0], validate True , time:0.07073379ms
-              out_f32x4_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.08693480ms
-              out_f32x4_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.02308369ms
-       out_f32x4_shared_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.02914500ms
-       out_f32x4_shared_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.03913331ms
-   out_f32x4_shared_bcf_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.02951288ms
-   out_f32x4_shared_bcf_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.03415990ms
- out_mat_transpose_cute_col2row_reg: [0.0, 2048.0, 1.0], validate True , time:0.05992270ms
- out_mat_transpose_cute_row2col_reg: [0.0, 2048.0, 1.0], validate True , time:0.04137969ms
-    out_mat_transpose_cute_col_smem: [0.0, 2048.0, 1.0], validate True , time:0.06176162ms
-    out_mat_transpose_cute_row_smem: [0.0, 2048.0, 1.0], validate True , time:0.04155755ms
-out_mat_transpose_cute_col_smem_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.05330944ms
-out_mat_transpose_cute_row_smem_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.03944159ms
-out_mat_transpose_cute_row_cvectorized: [0.0, 2048.0, 1.0], validate True , time:0.04380369ms
-out_mat_transpose_cute_row_rvectorized: [0.0, 2048.0, 1.0], validate True , time:0.02814078ms
-out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.03094649ms
-out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.03010845ms
-                         out_f32_th: [0.0, 2048.0, 1.0], validate True , time:0.09364438ms
-                out_f32_th_compiled: [0.0, 2048.0, 1.0], validate True , time:0.05807614ms
+                       out_original: [0.0, 1.0, 2048.0], validate False, time:0.08468652ms
+                    out_f32_col2row: [0.0, 2048.0, 1.0], validate True , time:0.25145221ms
+                    out_f32_row2col: [0.0, 2048.0, 1.0], validate True , time:0.16740918ms
+                out_f32_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.12791324ms
+                out_f32_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.12523937ms
+                  out_f32x4_col2row: [0.0, 2048.0, 1.0], validate True , time:0.26657438ms
+                  out_f32x4_row2col: [0.0, 2048.0, 1.0], validate True , time:0.15698409ms
+              out_f32x4_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.13683271ms
+              out_f32x4_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.08711052ms
+       out_f32x4_shared_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.08722329ms
+       out_f32x4_shared_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.08619714ms
+   out_f32x4_shared_bcf_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.08731914ms
+   out_f32x4_shared_bcf_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.08600688ms
+ out_mat_transpose_cute_col2row_reg: [0.0, 2048.0, 1.0], validate True , time:0.12273955ms
+ out_mat_transpose_cute_row2col_reg: [0.0, 2048.0, 1.0], validate True , time:0.12345028ms
+    out_mat_transpose_cute_col_smem: [0.0, 2048.0, 1.0], validate True , time:0.12603784ms
+    out_mat_transpose_cute_row_smem: [0.0, 2048.0, 1.0], validate True , time:0.12179136ms
+out_mat_transpose_cute_col_smem_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.12539601ms
+out_mat_transpose_cute_row_smem_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.12133408ms
+out_mat_transpose_cute_row_cvectorized: [0.0, 2048.0, 1.0], validate True , time:0.09775424ms
+out_mat_transpose_cute_row_rvectorized: [0.0, 2048.0, 1.0], validate True , time:0.08309197ms
+out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.09721851ms
+out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.08333635ms
+out_mat_transpose_cute_row_rvectorized_swizzled_optimized: [0.0, 2048.0, 1.0], validate True , time:0.08407569ms
+                         out_f32_th: [0.0, 2048.0, 1.0], validate True , time:0.24632931ms
+                out_f32_th_compiled: [0.0, 2048.0, 1.0], validate True , time:0.24452114ms
 ----------------------------------------------------------------------------------------------------------------------------------
 ----------------------------------------------------------------------------------------------------------------------------------
                                                        M=4096, N=4096
-                       out_original: [0.0, 1.0, 4096.0], validate False, time:0.04761362ms
-                    out_f32_col2row: [0.0, 4096.0, 1.0], validate True , time:0.26551914ms
-                    out_f32_row2col: [0.0, 4096.0, 1.0], validate True , time:0.12826180ms
-                out_f32_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.12965775ms
-                out_f32_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.12970257ms
-                  out_f32_diagnonal: [0.0, 4096.0, 1.0], validate True , time:0.08692741ms
-                  out_f32x4_col2row: [0.0, 4096.0, 1.0], validate True , time:0.27640176ms
-                  out_f32x4_row2col: [0.0, 4096.0, 1.0], validate True , time:0.13386226ms
-              out_f32x4_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.16657877ms
-              out_f32x4_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.04246473ms
-       out_f32x4_shared_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.05401349ms
-       out_f32x4_shared_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.07208681ms
-   out_f32x4_shared_bcf_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.05456352ms
-   out_f32x4_shared_bcf_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.06306529ms
- out_mat_transpose_cute_col2row_reg: [0.0, 4096.0, 1.0], validate True , time:0.11687827ms
- out_mat_transpose_cute_row2col_reg: [0.0, 4096.0, 1.0], validate True , time:0.07934999ms
-    out_mat_transpose_cute_col_smem: [0.0, 4096.0, 1.0], validate True , time:0.11962628ms
-    out_mat_transpose_cute_row_smem: [0.0, 4096.0, 1.0], validate True , time:0.07979536ms
-out_mat_transpose_cute_col_smem_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.10279226ms
-out_mat_transpose_cute_row_smem_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.07552862ms
-out_mat_transpose_cute_row_cvectorized: [0.0, 4096.0, 1.0], validate True , time:0.08200026ms
-out_mat_transpose_cute_row_rvectorized: [0.0, 4096.0, 1.0], validate True , time:0.05083799ms
-out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.05675173ms
-out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.05572939ms
-                         out_f32_th: [0.0, 4096.0, 1.0], validate True , time:0.17590189ms
-                out_f32_th_compiled: [0.0, 4096.0, 1.0], validate True , time:0.05817580ms
+                       out_original: [0.0, 1.0, 4096.0], validate False, time:0.16618252ms
+                    out_f32_col2row: [0.0, 4096.0, 1.0], validate True , time:0.50097990ms
+                    out_f32_row2col: [0.0, 4096.0, 1.0], validate True , time:0.35535216ms
+                out_f32_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.27060771ms
+                out_f32_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.25861597ms
+                  out_f32_diagnonal: [0.0, 4096.0, 1.0], validate True , time:0.26617074ms
+                  out_f32x4_col2row: [0.0, 4096.0, 1.0], validate True , time:0.51177859ms
+                  out_f32x4_row2col: [0.0, 4096.0, 1.0], validate True , time:0.31568432ms
+              out_f32x4_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.26652837ms
+              out_f32x4_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.17010331ms
+       out_f32x4_shared_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.17163038ms
+       out_f32x4_shared_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.16799879ms
+   out_f32x4_shared_bcf_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.17183828ms
+   out_f32x4_shared_bcf_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.16800833ms
+ out_mat_transpose_cute_col2row_reg: [0.0, 4096.0, 1.0], validate True , time:0.26187921ms
+ out_mat_transpose_cute_row2col_reg: [0.0, 4096.0, 1.0], validate True , time:0.24576879ms
+    out_mat_transpose_cute_col_smem: [0.0, 4096.0, 1.0], validate True , time:0.25347733ms
+    out_mat_transpose_cute_row_smem: [0.0, 4096.0, 1.0], validate True , time:0.24340343ms
+out_mat_transpose_cute_col_smem_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.25243449ms
+out_mat_transpose_cute_row_smem_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.24239349ms
+out_mat_transpose_cute_row_cvectorized: [0.0, 4096.0, 1.0], validate True , time:0.19701624ms
+out_mat_transpose_cute_row_rvectorized: [0.0, 4096.0, 1.0], validate True , time:0.16537905ms
+out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.19600892ms
+out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.16557527ms
+out_mat_transpose_cute_row_rvectorized_swizzled_optimized: [0.0, 4096.0, 1.0], validate True , time:0.16697907ms
+                         out_f32_th: [0.0, 4096.0, 1.0], validate True , time:0.49465632ms
+                out_f32_th_compiled: [0.0, 4096.0, 1.0], validate True , time:0.49463820ms
 ----------------------------------------------------------------------------------------------------------------------------------
 ----------------------------------------------------------------------------------------------------------------------------------
                                                        M=4096, N=8192
-                       out_original: [0.0, 1.0, 8192.0], validate False, time:0.09065676ms
-                    out_f32_col2row: [0.0, 8192.0, 1.0], validate True , time:0.52875924ms
-                    out_f32_row2col: [0.0, 8192.0, 1.0], validate True , time:0.24883223ms
-                out_f32_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.25138545ms
-                out_f32_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.25151587ms
-                  out_f32x4_col2row: [0.0, 8192.0, 1.0], validate True , time:0.55523229ms
-                  out_f32x4_row2col: [0.0, 8192.0, 1.0], validate True , time:0.25797868ms
-              out_f32x4_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.33464503ms
-              out_f32x4_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.08137870ms
-       out_f32x4_shared_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.10357594ms
-       out_f32x4_shared_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.13923311ms
-   out_f32x4_shared_bcf_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.10485888ms
-   out_f32x4_shared_bcf_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.12177515ms
- out_mat_transpose_cute_col2row_reg: [0.0, 8192.0, 1.0], validate True , time:0.22806001ms
- out_mat_transpose_cute_row2col_reg: [0.0, 8192.0, 1.0], validate True , time:0.15632129ms
-    out_mat_transpose_cute_col_smem: [0.0, 8192.0, 1.0], validate True , time:0.23541665ms
-    out_mat_transpose_cute_row_smem: [0.0, 8192.0, 1.0], validate True , time:0.15710807ms
-out_mat_transpose_cute_col_smem_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.20194674ms
-out_mat_transpose_cute_row_smem_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.14867640ms
-out_mat_transpose_cute_row_cvectorized: [0.0, 8192.0, 1.0], validate True , time:0.15857029ms
-out_mat_transpose_cute_row_rvectorized: [0.0, 8192.0, 1.0], validate True , time:0.09695339ms
-out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.10866976ms
-out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.10696220ms
-                         out_f32_th: [0.0, 8192.0, 1.0], validate True , time:0.33883190ms
-                out_f32_th_compiled: [0.0, 8192.0, 1.0], validate True , time:0.07038069ms
+                       out_original: [0.0, 1.0, 8192.0], validate False, time:0.32828760ms
+                    out_f32_col2row: [0.0, 8192.0, 1.0], validate True , time:1.02160501ms
+                    out_f32_row2col: [0.0, 8192.0, 1.0], validate True , time:0.72063184ms
+                out_f32_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.53131676ms
+                out_f32_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.52935743ms
+                  out_f32x4_col2row: [0.0, 8192.0, 1.0], validate True , time:1.03164721ms
+                  out_f32x4_row2col: [0.0, 8192.0, 1.0], validate True , time:0.80112934ms
+              out_f32x4_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.53537488ms
+              out_f32x4_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.34326696ms
+       out_f32x4_shared_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.35909820ms
+       out_f32x4_shared_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.33556294ms
+   out_f32x4_shared_bcf_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.35932803ms
+   out_f32x4_shared_bcf_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.33499861ms
+ out_mat_transpose_cute_col2row_reg: [0.0, 8192.0, 1.0], validate True , time:0.54176092ms
+ out_mat_transpose_cute_row2col_reg: [0.0, 8192.0, 1.0], validate True , time:0.48399639ms
+    out_mat_transpose_cute_col_smem: [0.0, 8192.0, 1.0], validate True , time:0.49951982ms
+    out_mat_transpose_cute_row_smem: [0.0, 8192.0, 1.0], validate True , time:0.48039985ms
+out_mat_transpose_cute_col_smem_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.49459934ms
+out_mat_transpose_cute_row_smem_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.47925377ms
+out_mat_transpose_cute_row_cvectorized: [0.0, 8192.0, 1.0], validate True , time:0.42132831ms
+out_mat_transpose_cute_row_rvectorized: [0.0, 8192.0, 1.0], validate True , time:0.32764554ms
+out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.41889501ms
+out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.32755017ms
+out_mat_transpose_cute_row_rvectorized_swizzled_optimized: [0.0, 8192.0, 1.0], validate True , time:0.33102536ms
+                         out_f32_th: [0.0, 8192.0, 1.0], validate True , time:0.99227643ms
+                out_f32_th_compiled: [0.0, 8192.0, 1.0], validate True , time:0.99217844ms
 ----------------------------------------------------------------------------------------------------------------------------------
 ----------------------------------------------------------------------------------------------------------------------------------
                                                        M=8192, N=1024
-                       out_original: [0.0, 1.0, 1024.0], validate False, time:0.02613044ms
-                    out_f32_col2row: [0.0, 1024.0, 1.0], validate True , time:0.14240122ms
-                    out_f32_row2col: [0.0, 1024.0, 1.0], validate True , time:0.06301308ms
-                out_f32_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.06837869ms
-                out_f32_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.06840038ms
-                  out_f32x4_col2row: [0.0, 1024.0, 1.0], validate True , time:0.14549279ms
-                  out_f32x4_row2col: [0.0, 1024.0, 1.0], validate True , time:0.06775117ms
-              out_f32x4_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.08498430ms
-              out_f32x4_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.02311444ms
-       out_f32x4_shared_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.02915359ms
-       out_f32x4_shared_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.03713298ms
-   out_f32x4_shared_bcf_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.02951527ms
-   out_f32x4_shared_bcf_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.03300023ms
- out_mat_transpose_cute_col2row_reg: [0.0, 1024.0, 1.0], validate True , time:0.05988598ms
- out_mat_transpose_cute_row2col_reg: [0.0, 1024.0, 1.0], validate True , time:0.04108763ms
-    out_mat_transpose_cute_col_smem: [0.0, 1024.0, 1.0], validate True , time:0.06168962ms
-    out_mat_transpose_cute_row_smem: [0.0, 1024.0, 1.0], validate True , time:0.04137659ms
-out_mat_transpose_cute_col_smem_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.05313158ms
-out_mat_transpose_cute_row_smem_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.03937101ms
-out_mat_transpose_cute_row_cvectorized: [0.0, 1024.0, 1.0], validate True , time:0.04361272ms
-out_mat_transpose_cute_row_rvectorized: [0.0, 1024.0, 1.0], validate True , time:0.02782607ms
-out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.03081965ms
-out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.02822685ms
-                         out_f32_th: [0.0, 1024.0, 1.0], validate True , time:0.08901119ms
-                out_f32_th_compiled: [0.0, 1024.0, 1.0], validate True , time:0.05782843ms
+                       out_original: [0.0, 1.0, 1024.0], validate False, time:0.08466172ms
+                    out_f32_col2row: [0.0, 1024.0, 1.0], validate True , time:0.25207758ms
+                    out_f32_row2col: [0.0, 1024.0, 1.0], validate True , time:0.18354273ms
+                out_f32_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.12594485ms
+                out_f32_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.12563300ms
+                  out_f32x4_col2row: [0.0, 1024.0, 1.0], validate True , time:0.27044630ms
+                  out_f32x4_row2col: [0.0, 1024.0, 1.0], validate True , time:0.16568422ms
+              out_f32x4_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.13587356ms
+              out_f32x4_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.08572435ms
+       out_f32x4_shared_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.08661747ms
+       out_f32x4_shared_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.08503509ms
+   out_f32x4_shared_bcf_col2row(2d): [0.0, 1024.0, 1.0], validate True , time:0.08664989ms
+   out_f32x4_shared_bcf_row2col(2d): [0.0, 1024.0, 1.0], validate True , time:0.08469391ms
+ out_mat_transpose_cute_col2row_reg: [0.0, 1024.0, 1.0], validate True , time:0.12477684ms
+ out_mat_transpose_cute_row2col_reg: [0.0, 1024.0, 1.0], validate True , time:0.11948586ms
+    out_mat_transpose_cute_col_smem: [0.0, 1024.0, 1.0], validate True , time:0.12046981ms
+    out_mat_transpose_cute_row_smem: [0.0, 1024.0, 1.0], validate True , time:0.11931205ms
+out_mat_transpose_cute_col_smem_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.11980653ms
+out_mat_transpose_cute_row_smem_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.11914992ms
+out_mat_transpose_cute_row_cvectorized: [0.0, 1024.0, 1.0], validate True , time:0.10827446ms
+out_mat_transpose_cute_row_rvectorized: [0.0, 1024.0, 1.0], validate True , time:0.08326149ms
+out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.10820341ms
+out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 1024.0, 1.0], validate True , time:0.08342290ms
+out_mat_transpose_cute_row_rvectorized_swizzled_optimized: [0.0, 1024.0, 1.0], validate True , time:0.08371043ms
+                         out_f32_th: [0.0, 1024.0, 1.0], validate True , time:0.27149653ms
+                out_f32_th_compiled: [0.0, 1024.0, 1.0], validate True , time:0.27161312ms
 ----------------------------------------------------------------------------------------------------------------------------------
 ----------------------------------------------------------------------------------------------------------------------------------
                                                        M=8192, N=2048
-                       out_original: [0.0, 1.0, 2048.0], validate False, time:0.04759407ms
-                    out_f32_col2row: [0.0, 2048.0, 1.0], validate True , time:0.27500105ms
-                    out_f32_row2col: [0.0, 2048.0, 1.0], validate True , time:0.12144923ms
-                out_f32_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.13192320ms
-                out_f32_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.13188601ms
-                  out_f32x4_col2row: [0.0, 2048.0, 1.0], validate True , time:0.29046535ms
-                  out_f32x4_row2col: [0.0, 2048.0, 1.0], validate True , time:0.13048077ms
-              out_f32x4_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.17475128ms
-              out_f32x4_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.04228449ms
-       out_f32x4_shared_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.05392790ms
-       out_f32x4_shared_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.06848621ms
-   out_f32x4_shared_bcf_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.05463552ms
-   out_f32x4_shared_bcf_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.06046867ms
- out_mat_transpose_cute_col2row_reg: [0.0, 2048.0, 1.0], validate True , time:0.11590195ms
- out_mat_transpose_cute_row2col_reg: [0.0, 2048.0, 1.0], validate True , time:0.07939339ms
-    out_mat_transpose_cute_col_smem: [0.0, 2048.0, 1.0], validate True , time:0.11953950ms
-    out_mat_transpose_cute_row_smem: [0.0, 2048.0, 1.0], validate True , time:0.07990146ms
-out_mat_transpose_cute_col_smem_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.10252953ms
-out_mat_transpose_cute_row_smem_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.07594895ms
-out_mat_transpose_cute_row_cvectorized: [0.0, 2048.0, 1.0], validate True , time:0.08208942ms
-out_mat_transpose_cute_row_rvectorized: [0.0, 2048.0, 1.0], validate True , time:0.05093598ms
-out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.05676866ms
-out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.05595016ms
-                         out_f32_th: [0.0, 2048.0, 1.0], validate True , time:0.17065334ms
-                out_f32_th_compiled: [0.0, 2048.0, 1.0], validate True , time:0.05812144ms
+                       out_original: [0.0, 1.0, 2048.0], validate False, time:0.16627026ms
+                    out_f32_col2row: [0.0, 2048.0, 1.0], validate True , time:0.50467086ms
+                    out_f32_row2col: [0.0, 2048.0, 1.0], validate True , time:0.39958572ms
+                out_f32_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.26184535ms
+                out_f32_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.25405502ms
+                  out_f32x4_col2row: [0.0, 2048.0, 1.0], validate True , time:0.51482177ms
+                  out_f32x4_row2col: [0.0, 2048.0, 1.0], validate True , time:0.35543489ms
+              out_f32x4_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.26602769ms
+              out_f32x4_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.16936231ms
+       out_f32x4_shared_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.17157674ms
+       out_f32x4_shared_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.16717434ms
+   out_f32x4_shared_bcf_col2row(2d): [0.0, 2048.0, 1.0], validate True , time:0.17186403ms
+   out_f32x4_shared_bcf_row2col(2d): [0.0, 2048.0, 1.0], validate True , time:0.16655135ms
+ out_mat_transpose_cute_col2row_reg: [0.0, 2048.0, 1.0], validate True , time:0.24904299ms
+ out_mat_transpose_cute_row2col_reg: [0.0, 2048.0, 1.0], validate True , time:0.24884462ms
+    out_mat_transpose_cute_col_smem: [0.0, 2048.0, 1.0], validate True , time:0.25050330ms
+    out_mat_transpose_cute_row_smem: [0.0, 2048.0, 1.0], validate True , time:0.24837542ms
+out_mat_transpose_cute_col_smem_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.25063992ms
+out_mat_transpose_cute_row_smem_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.24788022ms
+out_mat_transpose_cute_row_cvectorized: [0.0, 2048.0, 1.0], validate True , time:0.22211933ms
+out_mat_transpose_cute_row_rvectorized: [0.0, 2048.0, 1.0], validate True , time:0.16454816ms
+out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.22272992ms
+out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 2048.0, 1.0], validate True , time:0.16480446ms
+out_mat_transpose_cute_row_rvectorized_swizzled_optimized: [0.0, 2048.0, 1.0], validate True , time:0.16556239ms
+                         out_f32_th: [0.0, 2048.0, 1.0], validate True , time:0.56485367ms
+                out_f32_th_compiled: [0.0, 2048.0, 1.0], validate True , time:0.56441379ms
 ----------------------------------------------------------------------------------------------------------------------------------
 ----------------------------------------------------------------------------------------------------------------------------------
                                                        M=8192, N=4096
-                       out_original: [0.0, 1.0, 4096.0], validate False, time:0.09060454ms
-                    out_f32_col2row: [0.0, 4096.0, 1.0], validate True , time:0.53396153ms
-                    out_f32_row2col: [0.0, 4096.0, 1.0], validate True , time:0.23736954ms
-                out_f32_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.25855660ms
-                out_f32_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.25850701ms
-                  out_f32x4_col2row: [0.0, 4096.0, 1.0], validate True , time:0.56433153ms
-                  out_f32x4_row2col: [0.0, 4096.0, 1.0], validate True , time:0.25520277ms
-              out_f32x4_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.34376645ms
-              out_f32x4_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.08140516ms
-       out_f32x4_shared_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.10327840ms
-       out_f32x4_shared_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.13143134ms
-   out_f32x4_shared_bcf_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.10436797ms
-   out_f32x4_shared_bcf_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.11549497ms
- out_mat_transpose_cute_col2row_reg: [0.0, 4096.0, 1.0], validate True , time:0.22957921ms
- out_mat_transpose_cute_row2col_reg: [0.0, 4096.0, 1.0], validate True , time:0.15648627ms
-    out_mat_transpose_cute_col_smem: [0.0, 4096.0, 1.0], validate True , time:0.23523784ms
-    out_mat_transpose_cute_row_smem: [0.0, 4096.0, 1.0], validate True , time:0.15738535ms
-out_mat_transpose_cute_col_smem_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.20139265ms
-out_mat_transpose_cute_row_smem_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.14952707ms
-out_mat_transpose_cute_row_cvectorized: [0.0, 4096.0, 1.0], validate True , time:0.15815616ms
-out_mat_transpose_cute_row_rvectorized: [0.0, 4096.0, 1.0], validate True , time:0.09657621ms
-out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.10863233ms
-out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.10682678ms
-                         out_f32_th: [0.0, 4096.0, 1.0], validate True , time:0.33217621ms
-                out_f32_th_compiled: [0.0, 4096.0, 1.0], validate True , time:0.07057214ms
+                       out_original: [0.0, 1.0, 4096.0], validate False, time:0.32855296ms
+                    out_f32_col2row: [0.0, 4096.0, 1.0], validate True , time:1.02771664ms
+                    out_f32_row2col: [0.0, 4096.0, 1.0], validate True , time:0.82746363ms
+                out_f32_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.52390122ms
+                out_f32_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.53260994ms
+                  out_f32x4_col2row: [0.0, 4096.0, 1.0], validate True , time:1.02764344ms
+                  out_f32x4_row2col: [0.0, 4096.0, 1.0], validate True , time:0.73395181ms
+              out_f32x4_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.53176832ms
+              out_f32x4_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.33959198ms
+       out_f32x4_shared_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.35376525ms
+       out_f32x4_shared_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.33685756ms
+   out_f32x4_shared_bcf_col2row(2d): [0.0, 4096.0, 1.0], validate True , time:0.35392237ms
+   out_f32x4_shared_bcf_row2col(2d): [0.0, 4096.0, 1.0], validate True , time:0.33208728ms
+ out_mat_transpose_cute_col2row_reg: [0.0, 4096.0, 1.0], validate True , time:0.53561139ms
+ out_mat_transpose_cute_row2col_reg: [0.0, 4096.0, 1.0], validate True , time:0.50565004ms
+    out_mat_transpose_cute_col_smem: [0.0, 4096.0, 1.0], validate True , time:0.51027036ms
+    out_mat_transpose_cute_row_smem: [0.0, 4096.0, 1.0], validate True , time:0.50410295ms
+out_mat_transpose_cute_col_smem_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.50999451ms
+out_mat_transpose_cute_row_smem_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.50347900ms
+out_mat_transpose_cute_row_cvectorized: [0.0, 4096.0, 1.0], validate True , time:0.44723439ms
+out_mat_transpose_cute_row_rvectorized: [0.0, 4096.0, 1.0], validate True , time:0.33080840ms
+out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.44853663ms
+out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 4096.0, 1.0], validate True , time:0.33074450ms
+out_mat_transpose_cute_row_rvectorized_swizzled_optimized: [0.0, 4096.0, 1.0], validate True , time:0.33293080ms
+                         out_f32_th: [0.0, 4096.0, 1.0], validate True , time:1.16223931ms
+                out_f32_th_compiled: [0.0, 4096.0, 1.0], validate True , time:1.16229939ms
 ----------------------------------------------------------------------------------------------------------------------------------
 ----------------------------------------------------------------------------------------------------------------------------------
                                                        M=8192, N=8192
-                       out_original: [0.0, 1.0, 8192.0], validate False, time:0.17657781ms
-                    out_f32_col2row: [0.0, 8192.0, 1.0], validate True , time:1.05990601ms
-                    out_f32_row2col: [0.0, 8192.0, 1.0], validate True , time:0.46211433ms
-                out_f32_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.50515103ms
-                out_f32_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.50519037ms
-                  out_f32_diagnonal: [0.0, 8192.0, 1.0], validate True , time:0.34358025ms
-                  out_f32x4_col2row: [0.0, 8192.0, 1.0], validate True , time:1.10640693ms
-                  out_f32x4_row2col: [0.0, 8192.0, 1.0], validate True , time:0.49680948ms
-              out_f32x4_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.67155337ms
-              out_f32x4_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.15931034ms
-       out_f32x4_shared_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.20235109ms
-       out_f32x4_shared_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.25867033ms
-   out_f32x4_shared_bcf_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.20493627ms
-   out_f32x4_shared_bcf_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.22880530ms
- out_mat_transpose_cute_col2row_reg: [0.0, 8192.0, 1.0], validate True , time:0.45240569ms
- out_mat_transpose_cute_row2col_reg: [0.0, 8192.0, 1.0], validate True , time:0.31271243ms
-    out_mat_transpose_cute_col_smem: [0.0, 8192.0, 1.0], validate True , time:0.46673679ms
-    out_mat_transpose_cute_row_smem: [0.0, 8192.0, 1.0], validate True , time:0.31364536ms
-out_mat_transpose_cute_col_smem_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.39961338ms
-out_mat_transpose_cute_row_smem_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.29810309ms
-out_mat_transpose_cute_row_cvectorized: [0.0, 8192.0, 1.0], validate True , time:0.31140137ms
-out_mat_transpose_cute_row_rvectorized: [0.0, 8192.0, 1.0], validate True , time:0.18832517ms
-out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.21277142ms
-out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.20939684ms
-                         out_f32_th: [0.0, 8192.0, 1.0], validate True , time:0.64468265ms
-                out_f32_th_compiled: [0.0, 8192.0, 1.0], validate True , time:0.13783288ms
+                       out_original: [0.0, 1.0, 8192.0], validate False, time:0.65326476ms
+                    out_f32_col2row: [0.0, 8192.0, 1.0], validate True , time:2.06266975ms
+                    out_f32_row2col: [0.0, 8192.0, 1.0], validate True , time:1.63385463ms
+                out_f32_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:1.06545210ms
+                out_f32_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:1.06657267ms
+                  out_f32_diagnonal: [0.0, 8192.0, 1.0], validate True , time:1.11862946ms
+                  out_f32x4_col2row: [0.0, 8192.0, 1.0], validate True , time:2.06704783ms
+                  out_f32x4_row2col: [0.0, 8192.0, 1.0], validate True , time:1.67259908ms
+              out_f32x4_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:1.06197333ms
+              out_f32x4_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.69753051ms
+       out_f32x4_shared_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.75325727ms
+       out_f32x4_shared_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.66890717ms
+   out_f32x4_shared_bcf_col2row(2d): [0.0, 8192.0, 1.0], validate True , time:0.75367093ms
+   out_f32x4_shared_bcf_row2col(2d): [0.0, 8192.0, 1.0], validate True , time:0.66564393ms
+ out_mat_transpose_cute_col2row_reg: [0.0, 8192.0, 1.0], validate True , time:1.04156613ms
+ out_mat_transpose_cute_row2col_reg: [0.0, 8192.0, 1.0], validate True , time:0.98622823ms
+    out_mat_transpose_cute_col_smem: [0.0, 8192.0, 1.0], validate True , time:1.01331401ms
+    out_mat_transpose_cute_row_smem: [0.0, 8192.0, 1.0], validate True , time:0.98386550ms
+out_mat_transpose_cute_col_smem_swizzled: [0.0, 8192.0, 1.0], validate True , time:1.00588274ms
+out_mat_transpose_cute_row_smem_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.98301244ms
+out_mat_transpose_cute_row_cvectorized: [0.0, 8192.0, 1.0], validate True , time:0.90383577ms
+out_mat_transpose_cute_row_rvectorized: [0.0, 8192.0, 1.0], validate True , time:0.65374827ms
+out_mat_transpose_cute_row_cvectorized_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.90814710ms
+out_mat_transpose_cute_row_rvectorized_swizzled: [0.0, 8192.0, 1.0], validate True , time:0.65394568ms
+out_mat_transpose_cute_row_rvectorized_swizzled_optimized: [0.0, 8192.0, 1.0], validate True , time:0.65754938ms
+                         out_f32_th: [0.0, 8192.0, 1.0], validate True , time:2.29676986ms
+                out_f32_th_compiled: [0.0, 8192.0, 1.0], validate True , time:2.29685950ms
 ----------------------------------------------------------------------------------------------------------------------------------
 
 ```
diff --git a/kernels/mat-transpose/mat_transpose.cu b/kernels/mat-transpose/mat_transpose.cu
index e3a06815..68fde378 100644
--- a/kernels/mat-transpose/mat_transpose.cu
+++ b/kernels/mat-transpose/mat_transpose.cu
@@ -377,6 +377,9 @@ extern void mat_transpose_cute_row_cvectorized_swizzled(torch::Tensor,
                                                         torch::Tensor);
 extern void mat_transpose_cute_row_rvectorized_swizzled(torch::Tensor,
                                                         torch::Tensor);
+extern void
+    mat_transpose_cute_row_rvectorized_swizzled_optimized(torch::Tensor,
+                                                          torch::Tensor);
 
 PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
   // 1d index
@@ -408,4 +411,6 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
   TORCH_BINDING_COMMON_EXTENSION(mat_transpose_cute_row_rvectorized)
   TORCH_BINDING_COMMON_EXTENSION(mat_transpose_cute_row_cvectorized_swizzled)
   TORCH_BINDING_COMMON_EXTENSION(mat_transpose_cute_row_rvectorized_swizzled)
+  TORCH_BINDING_COMMON_EXTENSION(
+      mat_transpose_cute_row_rvectorized_swizzled_optimized)
 }
diff --git a/kernels/mat-transpose/mat_transpose.py b/kernels/mat-transpose/mat_transpose.py
index 0c5e80ab..75fcd45f 100644
--- a/kernels/mat-transpose/mat_transpose.py
+++ b/kernels/mat-transpose/mat_transpose.py
@@ -4,12 +4,15 @@
 from typing import Optional
 
 import torch
+import torch._dynamo
 from torch.utils.cpp_extension import load
 
+torch._dynamo.config.suppress_errors = True
+
 torch.set_grad_enabled(False)
 
 CUTLASS_REPO_PATH = os.environ.get(
-    "CUTLASS_REPO_PATH", os.path.expanduser("~/cutlass")
+    "CUTLASS_REPO_PATH", os.path.expanduser("../../third-party/cutlass")
 )
 # Load the CUDA kernel as a python module
 lib = load(
@@ -178,6 +181,12 @@ def transpose_copy_compiled(input: torch.Tensor, out: torch.Tensor):
         "mat_transpose_cute_row_rvectorized_swizzled",
         y,
     )
+    run_benchmark(
+        lib.mat_transpose_cute_row_rvectorized_swizzled_optimized,
+        x,
+        "mat_transpose_cute_row_rvectorized_swizzled_optimized",
+        y,
+    )
     run_benchmark(
         partial(torch.transpose_copy, dim0=0, dim1=1, out=y), x, "f32_th"
     )
diff --git a/kernels/mat-transpose/mat_transpose_cute.cu b/kernels/mat-transpose/mat_transpose_cute.cu
index 594d4604..2ff63772 100644
--- a/kernels/mat-transpose/mat_transpose_cute.cu
+++ b/kernels/mat-transpose/mat_transpose_cute.cu
@@ -440,3 +440,97 @@ void mat_transpose_cute_row_rvectorized_swizzled(torch::Tensor x,
       ptr_A, ptr_B, M, N, tile_copy_a, tile_copy_b, sA_layout, sB_layout);
   CUDA_CHECK(cudaGetLastError());
 }
+
+template <typename T, int BLK_M, int BLK_N, typename TiledCopyA,
+          typename TiledCopyTrans, typename TiledCopyB, typename SmemLayoutB>
+__global__ void mat_transpose_cute_smem_vectorized_optimized_kernel(
+    const T *pA, T *pB, int M, int N, TiledCopyA copy_a,
+    TiledCopyTrans copy_trans, TiledCopyB copy_b, SmemLayoutB sB_layout) {
+  int tx = threadIdx.x;
+  int bx = blockIdx.x, by = blockIdx.y;
+
+  auto mA = make_tensor(make_gmem_ptr(pA),
+                        make_layout(make_shape(M, N), GenRowMajor{})); // (M, N)
+  auto mB = make_tensor(make_gmem_ptr(pB),
+                        make_layout(make_shape(N, M), GenRowMajor{})); // (N, N)
+
+  auto gA = local_tile(mA, make_shape(Int<BLK_M>{}, Int<BLK_N>{}),
+                       make_coord(bx, by)); // (BM, BN)
+  auto gB = local_tile(mB, make_shape(Int<BLK_N>{}, Int<BLK_M>{}),
+                       make_coord(by, bx)); // (BN, BM)
+
+  __shared__ T smem[BLK_M * BLK_N];
+  auto sB = make_tensor(make_smem_ptr(smem),
+                        sB_layout); // (BN, BM)
+
+  auto thr_copy_a = copy_a.get_slice(tx);
+  Tensor tAgA = thr_copy_a.partition_S(gA);
+  auto tAsA = make_tensor_like(tAgA);
+  Tensor tAsA_view = thr_copy_a.retile_D(tAsA);
+  copy(copy_a, tAgA, tAsA_view);
+
+  auto thr_copy_trans = copy_trans.get_slice(tx);
+  auto tAsB = thr_copy_trans.retile_S(tAsA);
+  auto tBsB_trans = thr_copy_trans.partition_D(sB);
+  copy(copy_trans, tAsB, tBsB_trans);
+
+  auto thr_copy_b = copy_b.get_slice(tx);
+  Tensor tBsB = thr_copy_b.partition_S(sB);
+  Tensor tBgB = thr_copy_b.partition_D(gB);
+
+  copy(copy_b, tBsB, tBgB);
+}
+
+void mat_transpose_cute_row_rvectorized_swizzled_optimized(torch::Tensor x,
+                                                           torch::Tensor y) {
+  const int BM = 8;
+  const int BN = 16 * 8;
+  auto ptr_A = x.data_ptr<float>();
+  auto ptr_B = y.data_ptr<float>();
+  const int M = x.size(0);
+  const int N = x.size(1);
+
+  // sanity checks
+  assert(M % 4 == 0);
+  assert(N % 4 == 0);
+  static_assert(BM % 4 == 0);
+  static_assert(BN % 4 == 0);
+  assert(is_aligned_128(ptr_A));
+  assert(is_aligned_128(ptr_B));
+
+  // 一次性加载8*16大小的矩阵
+  auto tile_copy_a = make_tiled_copy(
+      Copy_Atom<AutoVectorizingCopy, float>{},
+      make_layout(make_shape(Int<BM>{}, make_shape(Int<4>{}, Int<BN / 16>{})),
+                  make_stride(Int<4>{}, make_stride(Int<1>{}, Int<32>{}))),
+      make_layout(make_shape(Int<1>{}, Int<4>{}), GenRowMajor{}));
+
+  // 转换数据
+  auto tile_copy_trans = make_tiled_copy(
+      Copy_Atom<AutoVectorizingCopy, float>{},
+      make_layout(make_shape(make_shape(Int<4>{}, Int<BN / 16>{}), Int<BM>{}),
+                  make_stride(make_stride(Int<1>{}, Int<32>{}), Int<4>{})),
+      make_layout(make_shape(Int<4>{}, Int<1>{}), GenRowMajor{}));
+
+  // 一次性存储16*8大小的矩阵
+  auto tile_copy_b = make_tiled_copy(
+      Copy_Atom<AutoVectorizingCopy, float>{},
+      make_layout(make_shape(Int<BN>{}, Int<BM / 4>{}), GenRowMajor{}),
+      make_layout(make_shape(Int<1>{}, Int<4>{}), GenRowMajor{}));
+
+  auto swizzle_func = Swizzle<2, 3, 2>{};
+  auto sB_layout =
+      composition(swizzle_func,
+                  make_layout(make_shape(Int<BN>{}, Int<BM>{}), GenRowMajor{}));
+
+  static_assert(size(tile_copy_a) == size(tile_copy_b));
+
+  dim3 block(size(tile_copy_a));
+  dim3 grid((M + BM - 1) / BM, (N + BN - 1) / BN);
+
+  mat_transpose_cute_smem_vectorized_optimized_kernel<
+      float, BM, BN, decltype(tile_copy_a), decltype(tile_copy_trans),
+      decltype(tile_copy_b), decltype(sB_layout)><<<grid, block>>>(
+      ptr_A, ptr_B, M, N, tile_copy_a, tile_copy_trans, tile_copy_b, sB_layout);
+  CUDA_CHECK(cudaGetLastError());
+}

From ace5f16e42c1dbf8e6a436289090165e21e61733 Mon Sep 17 00:00:00 2001
From: hebangwen <32662175+hebangwen@users.noreply.github.com>
Date: Tue, 10 Jun 2025 11:18:02 +0800
Subject: [PATCH 16/57] bugfix: fix layernorm & rmsnorm f16 overflow (#335)

* [Test] Add f16 overflow testcase in layernorm and rmsnorm

* [Fix] call correct `rms_norm_f16x8_f32` and fix epsilon position
---
 kernels/layer-norm/layer_norm.cu |  8 ++++----
 kernels/layer-norm/layer_norm.py | 12 ++++++++++++
 kernels/rms-norm/rms_norm.cu     | 16 ++++++++--------
 kernels/rms-norm/rms_norm.py     | 13 +++++++++++++
 4 files changed, 37 insertions(+), 12 deletions(-)

diff --git a/kernels/layer-norm/layer_norm.cu b/kernels/layer-norm/layer_norm.cu
index d86e4b00..6038d650 100644
--- a/kernels/layer-norm/layer_norm.cu
+++ b/kernels/layer-norm/layer_norm.cu
@@ -197,7 +197,7 @@ __global__ void layer_norm_f16_f16_kernel(half *x, half *y, float g, float b,
   half variance = (value - s_mean) * (value - s_mean);
   variance = block_reduce_sum_f16_f16<NUM_THREADS>(variance);
   if (tid == 0)
-    s_variance = hrsqrt(variance / (K_ + epsilon));
+    s_variance = hrsqrt(variance / K_ + epsilon);
   // wait for s_variance in shared memory to be ready for all threads
   __syncthreads();
   if (idx < N * K) {
@@ -232,7 +232,7 @@ __global__ void layer_norm_f16x2_f16_kernel(half *x, half *y, float g, float b,
   half variance = reg_x_hat.x * reg_x_hat.x + reg_x_hat.y * reg_x_hat.y;
   variance = block_reduce_sum_f16_f16<NUM_THREADS>(variance);
   if (tid == 0)
-    s_variance = hrsqrt(variance / (K_ + epsilon));
+    s_variance = hrsqrt(variance / K_ + epsilon);
   // wait for s_variance in shared memory to be ready for all threads
   __syncthreads();
   if (idx < N * K) {
@@ -300,7 +300,7 @@ __global__ void layer_norm_f16x8_f16_kernel(half *x, half *y, float g, float b,
 
   variance = block_reduce_sum_f16_f16<NUM_THREADS>(variance);
   if (tid == 0)
-    s_variance = hrsqrt(variance / (K_ + epsilon));
+    s_variance = hrsqrt(variance / K_ + epsilon);
   // wait for s_variance in shared memory to be ready for all threads
   __syncthreads();
   // manual unroll
@@ -390,7 +390,7 @@ __global__ void layer_norm_f16x8_pack_f16_kernel(half *x, half *y, float g,
   }
   variance = block_reduce_sum_f16_f16<NUM_THREADS>(variance);
   if (tid == 0)
-    s_variance = hrsqrt(variance / (K_ + epsilon));
+    s_variance = hrsqrt(variance / K_ + epsilon);
   // wait for s_variance in shared memory to be ready for all threads
   __syncthreads();
 
diff --git a/kernels/layer-norm/layer_norm.py b/kernels/layer-norm/layer_norm.py
index 380cda1e..e9c68400 100644
--- a/kernels/layer-norm/layer_norm.py
+++ b/kernels/layer-norm/layer_norm.py
@@ -96,6 +96,18 @@ def run_benchmark(
 run_benchmark(naive_layer_norm, x_f16, "f16_th")
 print("-" * 85)
 
+print(" " * 40 + f"f16 overflow without f32")
+print("-" * 85)
+x_f16 = x.half() * 100 # this will cause overflow for kernels without `f32`
+run_benchmark(lib.layer_norm_f16_f16, x_f16, "f16f16", out_f16)
+run_benchmark(lib.layer_norm_f16_f32, x_f16, "f16f32", out_f16)
+run_benchmark(lib.layer_norm_f16x2_f16, x_f16, "f16x2f16", out_f16)
+run_benchmark(lib.layer_norm_f16x8_f16, x_f16, "f16x8f16", out_f16)
+run_benchmark(lib.layer_norm_f16x8_pack_f16, x_f16, "f16x8packf16", out_f16)
+run_benchmark(lib.layer_norm_f16x8_pack_f32, x_f16, "f16x8packf32", out_f16)
+run_benchmark(naive_layer_norm, x_f16, "f16_th")
+print("-" * 85)
+
 print("-" * 85)
 N, K = 4096, 1024
 print(" " * 40 + f"N={N}, K={K}")
diff --git a/kernels/rms-norm/rms_norm.cu b/kernels/rms-norm/rms_norm.cu
index 11c3700b..1ade3303 100644
--- a/kernels/rms-norm/rms_norm.cu
+++ b/kernels/rms-norm/rms_norm.cu
@@ -172,7 +172,7 @@ __global__ void rms_norm_f16_f16_kernel(half *x, half *y, float g, int N,
   half variance = value * value;
   variance = block_reduce_sum_f16_f16<NUM_THREADS>(variance);
   if (tid == 0)
-    s_variance = hrsqrt(variance / (K_ + epsilon));
+    s_variance = hrsqrt(variance / K_ + epsilon);
   // wait for s_variance in shared memory to be ready for all threads
   __syncthreads();
   if (idx < N * K)
@@ -195,7 +195,7 @@ __global__ void rms_norm_f16x2_f16_kernel(half *x, half *y, float g, int N,
                                 : __float2half(0.0f);
   variance = block_reduce_sum_f16_f16<NUM_THREADS>(variance);
   if (tid == 0)
-    s_variance = hrsqrt(variance / (K_ + epsilon));
+    s_variance = hrsqrt(variance / K_ + epsilon);
   // wait for s_variance in shared memory to be ready for all threads
   __syncthreads();
   half2 reg_y;
@@ -241,7 +241,7 @@ __global__ void rms_norm_f16x8_f16_kernel(half *x, half *y, float g, int N,
   variance += HALF2_VARIANCE(reg_x_3, 6);
   variance = block_reduce_sum_f16_f16<NUM_THREADS>(variance);
   if (tid == 0)
-    s_variance = hrsqrt(variance / (K_ + epsilon));
+    s_variance = hrsqrt(variance / K_ + epsilon);
   // wait for s_variance in shared memory to be ready for all threads
   __syncthreads();
   // manual unroll
@@ -292,7 +292,7 @@ __global__ void rms_norm_f16x8_f32_kernel(half *x, half *y, float g, int N,
 
   variance = block_reduce_sum_f32<NUM_THREADS>(variance);
   if (tid == 0)
-    s_variance = rsqrtf(variance / ((float)K + epsilon));
+    s_variance = rsqrtf(variance / (float)K + epsilon);
   // wait for s_variance in shared memory to be ready for all threads
   __syncthreads();
   // manual unroll
@@ -328,7 +328,7 @@ __global__ void rms_norm_f16_f32_kernel(half *x, half *y, float g, int N,
   float variance = value * value;
   variance = block_reduce_sum_f32<NUM_THREADS>(variance);
   if (tid == 0)
-    s_variance = rsqrtf(variance / ((float)K + epsilon));
+    s_variance = rsqrtf(variance / (float)K + epsilon);
   // wait for s_variance in shared memory to be ready for all threads
   __syncthreads();
   if (idx < N * K) {
@@ -360,7 +360,7 @@ __global__ void rms_norm_f16x8_pack_f16_kernel(half *x, half *y, float g, int N,
   }
   variance = block_reduce_sum_f16_f16<NUM_THREADS>(variance);
   if (tid == 0)
-    s_variance = hrsqrt(variance / (K_ + epsilon));
+    s_variance = hrsqrt(variance / K_ + epsilon);
   // wait for s_variance in shared memory to be ready for all threads
   __syncthreads();
 
@@ -396,7 +396,7 @@ __global__ void rms_norm_f16x8_pack_f32_kernel(half *x, half *y, float g, int N,
   }
   variance = block_reduce_sum_f32<NUM_THREADS>(variance);
   if (tid == 0)
-    s_variance = rsqrtf(variance / ((float)K + epsilon));
+    s_variance = rsqrtf(variance / (float)K + epsilon);
   // wait for s_variance in shared memory to be ready for all threads
   __syncthreads();
 
@@ -626,7 +626,7 @@ void rms_norm_f32x4(torch::Tensor x, torch::Tensor y, float g) {
   }
 
 #define LANUCH_RMS_NORM_F16x8F32_KERNEL(K)                                     \
-  rms_norm_f16x8_f16_kernel<(K) / 8>                                           \
+  rms_norm_f16x8_f32_kernel<(K) / 8>                                           \
       <<<grid, block>>>(reinterpret_cast<half *>(x.data_ptr()),                \
                         reinterpret_cast<half *>(y.data_ptr()), g, N, (K));
 
diff --git a/kernels/rms-norm/rms_norm.py b/kernels/rms-norm/rms_norm.py
index 6823236a..913dcf38 100644
--- a/kernels/rms-norm/rms_norm.py
+++ b/kernels/rms-norm/rms_norm.py
@@ -95,6 +95,19 @@ def run_benchmark(
 run_benchmark(naive_rms_norm, x_f16, "f16_th")
 print("-" * 85)
 
+print(" " * 40 + f"f16 overflow without f32")
+print("-" * 85)
+x_f16 = x.half() * 100  # this will cause overflow for kernels without `f32`
+run_benchmark(lib.rms_norm_f16_f16, x_f16, "f16f16", out_f16)
+run_benchmark(lib.rms_norm_f16_f32, x_f16, "f16f32", out_f16)
+run_benchmark(lib.rms_norm_f16x2_f16, x_f16, "f16x2f16", out_f16)
+run_benchmark(lib.rms_norm_f16x8_f16, x_f16, "f16x8f16", out_f16)
+run_benchmark(lib.rms_norm_f16x8_f32, x_f16, "f16x8f32", out_f16)
+run_benchmark(lib.rms_norm_f16x8_pack_f16, x_f16, "f16x8packf16", out_f16)
+run_benchmark(lib.rms_norm_f16x8_pack_f32, x_f16, "f16x8packf32", out_f16)
+run_benchmark(naive_rms_norm, x_f16, "f16_th")
+print("-" * 85)
+
 print("-" * 85)
 N, K = 4096, 1024
 print(" " * 40 + f"N={N}, K={K}")

From 3bef2072e415e8bce16ba0d4647e753b0568a1a1 Mon Sep 17 00:00:00 2001
From: www <radioheads@163.com>
Date: Tue, 10 Jun 2025 23:46:59 +0800
Subject: [PATCH 17/57] Bugfix: fix a compilation error (#336)

< < < should be <<<
---
 kernels/mat-transpose/mat_transpose.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/kernels/mat-transpose/mat_transpose.cu b/kernels/mat-transpose/mat_transpose.cu
index 68fde378..318bc441 100644
--- a/kernels/mat-transpose/mat_transpose.cu
+++ b/kernels/mat-transpose/mat_transpose.cu
@@ -334,7 +334,7 @@ __global__ void mat_transpose_f32x4_shared_bcf_row2col2d_kernel(float *x,
     dim3 block(WARP_SIZE_S, WARP_SIZE_S);                                      \
     dim3 grid((N + WARP_SIZE_S - 1) / (WARP_SIZE_S * n_element_col),           \
               (M + WARP_SIZE_S - 1) / (WARP_SIZE_S * n_element_row));          \
-    mat_transpose_##tag##2d_kernel < < < grid,                                 \
+    mat_transpose_##tag##2d_kernel <<< grid,                                   \
         block >>> (reinterpret_cast<element_type *>(x.data_ptr()),             \
                    reinterpret_cast<element_type *>(y.data_ptr()), M, N);      \
   }

From 285310e975a733c87fb135800e8f270f4a4f14dd Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Mon, 16 Jun 2025 19:46:47 +0800
Subject: [PATCH 18/57] feat: Add DBCache release news (#338)

---
 README.md | 6 ++++++
 1 file changed, 6 insertions(+)

diff --git a/README.md b/README.md
index 87e40d20..3715002c 100644
--- a/README.md
+++ b/README.md
@@ -50,6 +50,12 @@
 ## 📖 News 🔥🔥
 <div id="news"></div>
 
+- [2025-06-16]: [DBCache](https://github.com/vipshop/DBCache) is release! A **Training-free UNet-style** Cache Acceleration for Diffusion Transformers! Take a try!  
+
+<div align='center'>
+  <img src='https://github.com/user-attachments/assets/a5ec4320-d2f9-4254-888a-170b2d9e3784' height="170px">
+</div>
+
 - [2025-01-08]: [📚Split Q + Fully QKV Fine-grained Tiling](#mma-tiling-qkv) has been refactored into 🤖[ffpa-attn](https://github.com/xlite-dev/ffpa-attn.git): 📚FFPA - Yet another Faster Flash Prefill Attention with O(1)🎉SRAM complexity for headdim > 256, **1.8x~3x**🎉faster than SDPA EA: [📈L20 ~1.9x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-l20), [📈 A30 ~1.8x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-a30), [📈3080 ~2.9x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-3080), [📈4090 ~2.1x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-4090).
 
 <div align='center'>

From 8b25c294604c56dfd59d2e022a26f20b551070e4 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Mon, 16 Jun 2025 19:48:17 +0800
Subject: [PATCH 19/57] Update README.md

---
 README.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/README.md b/README.md
index 3715002c..66886937 100644
--- a/README.md
+++ b/README.md
@@ -50,7 +50,7 @@
 ## 📖 News 🔥🔥
 <div id="news"></div>
 
-- [2025-06-16]: [DBCache](https://github.com/vipshop/DBCache) is release! A **Training-free UNet-style** Cache Acceleration for Diffusion Transformers! Take a try!  
+- [2025-06-16]: [DBCache: Dual Block Caching](https://github.com/vipshop/DBCache) is release! A **Training-free UNet-style** Cache Acceleration for **Diffusion Transformers**! Feel free to take a try!  
 
 <div align='center'>
   <img src='https://github.com/user-attachments/assets/a5ec4320-d2f9-4254-888a-170b2d9e3784' height="170px">

From 0cf539d3d65343fca3879078a51680a25b5ab1ff Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Mon, 16 Jun 2025 22:57:36 +0800
Subject: [PATCH 20/57] Update README.md

---
 README.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/README.md b/README.md
index 66886937..24bfe1ad 100644
--- a/README.md
+++ b/README.md
@@ -53,7 +53,7 @@
 - [2025-06-16]: [DBCache: Dual Block Caching](https://github.com/vipshop/DBCache) is release! A **Training-free UNet-style** Cache Acceleration for **Diffusion Transformers**! Feel free to take a try!  
 
 <div align='center'>
-  <img src='https://github.com/user-attachments/assets/a5ec4320-d2f9-4254-888a-170b2d9e3784' height="170px">
+  <img src='https://github.com/user-attachments/assets/a5ec4320-d2f9-4254-888a-170b2d9e3784' height=200px>
 </div>
 
 - [2025-01-08]: [📚Split Q + Fully QKV Fine-grained Tiling](#mma-tiling-qkv) has been refactored into 🤖[ffpa-attn](https://github.com/xlite-dev/ffpa-attn.git): 📚FFPA - Yet another Faster Flash Prefill Attention with O(1)🎉SRAM complexity for headdim > 256, **1.8x~3x**🎉faster than SDPA EA: [📈L20 ~1.9x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-l20), [📈 A30 ~1.8x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-a30), [📈3080 ~2.9x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-3080), [📈4090 ~2.1x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-4090).

From 6953b390155b58f35ccfa40264ce452ebd07ae16 Mon Sep 17 00:00:00 2001
From: Beatlesso <96542402+Beatlesso@users.noreply.github.com>
Date: Tue, 17 Jun 2025 00:27:27 +0800
Subject: [PATCH 21/57] add out_f32x4_shared_bcf_merge_write_row2col(2d) (#339)

---
 kernels/mat-transpose/mat_transpose.cu | 43 ++++++++++++++++++++++++++
 kernels/mat-transpose/mat_transpose.py |  6 ++++
 2 files changed, 49 insertions(+)

diff --git a/kernels/mat-transpose/mat_transpose.cu b/kernels/mat-transpose/mat_transpose.cu
index 318bc441..190115c3 100644
--- a/kernels/mat-transpose/mat_transpose.cu
+++ b/kernels/mat-transpose/mat_transpose.cu
@@ -216,6 +216,7 @@ __global__ void mat_transpose_f32x4_shared_row2col2d_kernel(float *x, float *y,
   }
 }
 
+
 __global__ void mat_transpose_f32x4_shared_bcf_col2row2d_kernel(float *x,
                                                                 float *y,
                                                                 const int row,
@@ -296,6 +297,44 @@ __global__ void mat_transpose_f32x4_shared_bcf_row2col2d_kernel(float *x,
     y[(out_y + 3) * row + out_x] = smem_val.w;
   }
 }
+
+
+__global__ void mat_transpose_f32x4_shared_bcf_merge_write_row2col2d_kernel(float *x,
+                                                                float *y,
+                                                                const int row,
+                                                                const int col) {
+  const int global_x = blockIdx.x * blockDim.x + threadIdx.x;
+  const int global_y = blockIdx.y * blockDim.y + threadIdx.y;
+  const int local_x = threadIdx.x;
+  const int local_y = threadIdx.y;
+  __shared__ float tile[WARP_SIZE_S * 4][WARP_SIZE_S + PAD];
+  if (global_y * 4 < row && global_x < col) {
+    // load value from x to shared memory
+    float4 x_val;
+    x_val.x = x[(global_y * 4) * col + global_x];
+    x_val.y = x[(global_y * 4 + 1) * col + global_x];
+    x_val.z = x[(global_y * 4 + 2) * col + global_x];
+    x_val.w = x[(global_y * 4 + 3) * col + global_x];
+    tile[local_y * 4][local_x] = x_val.x;
+    tile[local_y * 4 + 1][local_x] = x_val.y;
+    tile[local_y * 4 + 2][local_x] = x_val.z;
+    tile[local_y * 4 + 3][local_x] = x_val.w;
+    __syncthreads();
+    float4 smem_val;
+    // load value from shared memory to y.
+    smem_val.x = tile[local_x * 4][local_y];
+    smem_val.y = tile[local_x * 4 + 1][local_y];
+    smem_val.z = tile[local_x * 4 + 2][local_y];
+    smem_val.w = tile[local_x * 4 + 3][local_y];
+
+    const int gid_x = blockIdx.x * blockDim.x;
+    const int gid_y = blockIdx.y * blockDim.y * 4;    
+    const int out_y = gid_y + local_x * 4;
+    const int out_x = gid_x + local_y;
+    reinterpret_cast<float4 *>(y)[(out_x * row + out_y) / 4] = FLOAT4(smem_val);
+  }
+}
+
 // TODO: may support double buffer pipeline mat transpose ?
 // TODO: may support fp16 mat transpose ?
 
@@ -361,6 +400,9 @@ TORCH_BINDING_MAT_TRANSPOSE2D(f32x4_shared_bcf_col2row, torch::kFloat32, float,
                               1, 4)
 TORCH_BINDING_MAT_TRANSPOSE2D(f32x4_shared_bcf_row2col, torch::kFloat32, float,
                               4, 1)
+TORCH_BINDING_MAT_TRANSPOSE2D(f32x4_shared_bcf_merge_write_row2col, torch::kFloat32, float,
+                              4, 1)
+
 // TODO: may support double buffer pipeline mat transpose ?
 // TODO: may support fp16 mat transpose ?
 
@@ -400,6 +442,7 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
   // shared memory optimize with bcf
   TORCH_BINDING_COMMON_EXTENSION(mat_transpose_f32x4_shared_bcf_col2row2d)
   TORCH_BINDING_COMMON_EXTENSION(mat_transpose_f32x4_shared_bcf_row2col2d)
+  TORCH_BINDING_COMMON_EXTENSION(mat_transpose_f32x4_shared_bcf_merge_write_row2col2d)
   // CuTe implentations
   TORCH_BINDING_COMMON_EXTENSION(mat_transpose_cute_col2row_reg)
   TORCH_BINDING_COMMON_EXTENSION(mat_transpose_cute_row2col_reg)
diff --git a/kernels/mat-transpose/mat_transpose.py b/kernels/mat-transpose/mat_transpose.py
index 75fcd45f..62945183 100644
--- a/kernels/mat-transpose/mat_transpose.py
+++ b/kernels/mat-transpose/mat_transpose.py
@@ -127,6 +127,12 @@ def transpose_copy_compiled(input: torch.Tensor, out: torch.Tensor):
         "f32x4_shared_bcf_row2col(2d)",
         y,
     )
+    run_benchmark(
+        lib.mat_transpose_f32x4_shared_bcf_merge_write_row2col2d,
+        x,
+        "f32x4_shared_bcf_merge_write_row2col(2d)",
+        y,
+    )
     run_benchmark(
         lib.mat_transpose_cute_col2row_reg,
         x,

From fa3c456e86646b5d2795d530d21ffeb4098260fe Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Tue, 17 Jun 2025 19:56:20 +0800
Subject: [PATCH 22/57] Update README.md

---
 README.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/README.md b/README.md
index 24bfe1ad..edf8c375 100644
--- a/README.md
+++ b/README.md
@@ -50,7 +50,7 @@
 ## 📖 News 🔥🔥
 <div id="news"></div>
 
-- [2025-06-16]: [DBCache: Dual Block Caching](https://github.com/vipshop/DBCache) is release! A **Training-free UNet-style** Cache Acceleration for **Diffusion Transformers**! Feel free to take a try!  
+- [2025-06-16]: [🤗CacheDiT](https://github.com/vipshop/cache-dit) is release! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for DiT (**DBCache**, **DBPrune**, **FBCache**, etc). Feel free to take a try!  
 
 <div align='center'>
   <img src='https://github.com/user-attachments/assets/a5ec4320-d2f9-4254-888a-170b2d9e3784' height=200px>

From e2b490d89f45f4cf98ee0d1fa3bb58caa781e4d1 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Wed, 18 Jun 2025 08:53:06 +0800
Subject: [PATCH 23/57] Update README.md

---
 README.md | 14 +++++++-------
 1 file changed, 7 insertions(+), 7 deletions(-)

diff --git a/README.md b/README.md
index edf8c375..9d4ab4ae 100644
--- a/README.md
+++ b/README.md
@@ -53,23 +53,23 @@
 - [2025-06-16]: [🤗CacheDiT](https://github.com/vipshop/cache-dit) is release! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for DiT (**DBCache**, **DBPrune**, **FBCache**, etc). Feel free to take a try!  
 
 <div align='center'>
-  <img src='https://github.com/user-attachments/assets/a5ec4320-d2f9-4254-888a-170b2d9e3784' height=200px>
+  <img src='https://github.com/user-attachments/assets/a5ec4320-d2f9-4254-888a-170b2d9e3784' height=170px>
 </div>
 
 - [2025-01-08]: [📚Split Q + Fully QKV Fine-grained Tiling](#mma-tiling-qkv) has been refactored into 🤖[ffpa-attn](https://github.com/xlite-dev/ffpa-attn.git): 📚FFPA - Yet another Faster Flash Prefill Attention with O(1)🎉SRAM complexity for headdim > 256, **1.8x~3x**🎉faster than SDPA EA: [📈L20 ~1.9x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-l20), [📈 A30 ~1.8x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-a30), [📈3080 ~2.9x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-3080), [📈4090 ~2.1x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-4090).
 
 <div align='center'>
-  <img src='https://github.com/user-attachments/assets/cba2edce-ac0d-412e-823c-7eea2cc63f83' height="170px" width="270px">
-  <img src='https://github.com/user-attachments/assets/447e2937-f7c8-47c8-8550-8c0c71b910e6' height="170px" width="270px">
-  <img src='https://github.com/user-attachments/assets/65a8d564-8fa7-4d66-86b9-e238feb86143' height="170px" width="270px">
+  <img src='https://github.com/user-attachments/assets/cba2edce-ac0d-412e-823c-7eea2cc63f83' height="170px" width="229px">
+  <img src='https://github.com/user-attachments/assets/447e2937-f7c8-47c8-8550-8c0c71b910e6' height="170px" width="229px">
+  <img src='https://github.com/user-attachments/assets/65a8d564-8fa7-4d66-86b9-e238feb86143' height="170px" width="229px">
 </div>
 
 - [2024-12-02]: HGEMM MMA kernels has been refactored into 🤖[HGEMM](https://github.com/xlite-dev/HGEMM.git): ⚡️Write HGEMM from scratch using Tensor Cores with WMMA, MMA and CuTe API, achieve peak⚡️ performance.
 
 <div align='center'>
-  <img src='https://github.com/user-attachments/assets/71927ac9-72b3-4ce9-b0e2-788b5885bc99' height="170px" width="270px">
-  <img src='https://github.com/user-attachments/assets/05ef4f5e-d999-48ea-b58e-782cffb24e85' height="170px" width="270px">
-  <img src='https://github.com/user-attachments/assets/9472e970-c083-4b31-9252-3eeecc761078' height="170px" width="270px">
+  <img src='https://github.com/user-attachments/assets/71927ac9-72b3-4ce9-b0e2-788b5885bc99' height="170px" width="229px">
+  <img src='https://github.com/user-attachments/assets/05ef4f5e-d999-48ea-b58e-782cffb24e85' height="170px" width="229px">
+  <img src='https://github.com/user-attachments/assets/9472e970-c083-4b31-9252-3eeecc761078' height="170px" width="229px">
 </div>
 
 

From b29acdaf481730318c1e90f1195bce0e34b9925f Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Wed, 18 Jun 2025 16:54:37 +0800
Subject: [PATCH 24/57] Update README.md

---
 README.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/README.md b/README.md
index 9d4ab4ae..835ac265 100644
--- a/README.md
+++ b/README.md
@@ -50,7 +50,7 @@
 ## 📖 News 🔥🔥
 <div id="news"></div>
 
-- [2025-06-16]: [🤗CacheDiT](https://github.com/vipshop/cache-dit) is release! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for DiT (**DBCache**, **DBPrune**, **FBCache**, etc). Feel free to take a try!  
+- [2025-06-16]: [🤗CacheDiT](https://github.com/vipshop/cache-dit) is release! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for Diffusion Transformers (**DBCache**, **DBPrune**, **FBCache**, etc). Feel free to take a try!  
 
 <div align='center'>
   <img src='https://github.com/user-attachments/assets/a5ec4320-d2f9-4254-888a-170b2d9e3784' height=170px>

From a035056d97380dc929a22f43a3b93661c4847ee0 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Fri, 20 Jun 2025 18:36:24 +0800
Subject: [PATCH 25/57] Update README.md

---
 README.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/README.md b/README.md
index 835ac265..0d031c6e 100644
--- a/README.md
+++ b/README.md
@@ -50,7 +50,7 @@
 ## 📖 News 🔥🔥
 <div id="news"></div>
 
-- [2025-06-16]: [🤗CacheDiT](https://github.com/vipshop/cache-dit) is release! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for Diffusion Transformers (**DBCache**, **DBPrune**, **FBCache**, etc). Feel free to take a try!  
+- [2025-06-16]: [🤗CacheDiT](https://github.com/vipshop/cache-dit) is release! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for Diffusion Transformers (**DBCache**, **DBPrune**, **FBCache**, etc)🔥. Feel free to take a try!  
 
 <div align='center'>
   <img src='https://github.com/user-attachments/assets/a5ec4320-d2f9-4254-888a-170b2d9e3784' height=170px>

From 404a3e1c80ab1cb9bc602fac33db44ad3c5fdb77 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Sat, 21 Jun 2025 18:09:46 +0800
Subject: [PATCH 26/57] Update README.md

---
 README.md | 1 +
 1 file changed, 1 insertion(+)

diff --git a/README.md b/README.md
index 0d031c6e..225ac381 100644
--- a/README.md
+++ b/README.md
@@ -519,6 +519,7 @@ The kernels listed here will guide you through a step-by-step progression, rangi
 
 |📖 类型-标题|📖 作者| 📖 推荐 |
 |:---|:---|:---|
+| [[Diffusion推理]📖DiT推理加速综述: Caching](https://zhuanlan.zhihu.com/p/711223667)|@DefTruth|⭐️⭐️⭐|  
 | [[Triton编程][基础]📖Triton极简入门: Triton Vector Add](https://zhuanlan.zhihu.com/p/1902778199261291694)|@DefTruth|⭐️⭐️⭐|
 | [[Triton编程][基础]📖Triton Fused Softmax Kernel详解: 从Python源码到PTX](https://zhuanlan.zhihu.com/p/1899562146477609112)|@DefTruth|⭐️⭐️⭐|
 | [[Triton编程][基础]📖vLLM Triton Merge Attention States Kernel详解](https://zhuanlan.zhihu.com/p/1904937907703243110)|@DefTruth|⭐️⭐️⭐|

From dacab18030b7fbefd4f9be27cebbc80e49aa2137 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Fri, 27 Jun 2025 11:59:52 +0800
Subject: [PATCH 27/57] pin HelloGitHub (#342)

---
 README.md | 4 +++-
 1 file changed, 3 insertions(+), 1 deletion(-)

diff --git a/README.md b/README.md
index 225ac381..f68c8115 100644
--- a/README.md
+++ b/README.md
@@ -9,10 +9,12 @@
 </div>
 --->
 
+
 <div align="center">
   <p align="center">
     <h2>📚 LeetCUDA: Modern CUDA Learn Notes with PyTorch for Beginners 🐑</h2>
-    <img src='https://github.com/user-attachments/assets/b2578723-b7a7-4d8f-bcd1-5008947b808a' >
+    <img src='https://github.com/user-attachments/assets/b2578723-b7a7-4d8f-bcd1-5008947b808a' width="430" height="56" >
+    <a href="https://hellogithub.com/repository/98348655a96640ca8ddcbc298edc901d" target="_blank"><img src="https://api.hellogithub.com/v1/widgets/recommend.svg?rid=98348655a96640ca8ddcbc298edc901d&claim_uid=ofSCbzTmdeQk3FD&theme=dark" alt="Featured｜HelloGitHub" style="width: 250px; height: 54px;" width="250" height="54" /></a>
   </p>
   <div align='center'>
       <img src=https://cdn.rawgit.com/sindresorhus/awesome/d7305f38d29fed78fa85652e3a63e154dd8e8829/media/badge.svg >

From 98284eb2e74f15e1f3a6057e4509ac3cc6728d1d Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Fri, 27 Jun 2025 16:20:25 +0800
Subject: [PATCH 28/57] Update README.md (#343)

---
 README.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/README.md b/README.md
index f68c8115..f0813112 100644
--- a/README.md
+++ b/README.md
@@ -13,7 +13,7 @@
 <div align="center">
   <p align="center">
     <h2>📚 LeetCUDA: Modern CUDA Learn Notes with PyTorch for Beginners 🐑</h2>
-    <img src='https://github.com/user-attachments/assets/b2578723-b7a7-4d8f-bcd1-5008947b808a' width="430" height="56" >
+    <img src='https://github.com/user-attachments/assets/b2578723-b7a7-4d8f-bcd1-5008947b808a' width="380" height="56" >
     <a href="https://hellogithub.com/repository/98348655a96640ca8ddcbc298edc901d" target="_blank"><img src="https://api.hellogithub.com/v1/widgets/recommend.svg?rid=98348655a96640ca8ddcbc298edc901d&claim_uid=ofSCbzTmdeQk3FD&theme=dark" alt="Featured｜HelloGitHub" style="width: 250px; height: 54px;" width="250" height="54" /></a>
   </p>
   <div align='center'>

From 955557bc8c512c9511de8230a3954e0916944a3a Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Fri, 27 Jun 2025 16:56:13 +0800
Subject: [PATCH 29/57] Update README.md (#344)

---
 README.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/README.md b/README.md
index f0813112..6d307311 100644
--- a/README.md
+++ b/README.md
@@ -13,7 +13,7 @@
 <div align="center">
   <p align="center">
     <h2>📚 LeetCUDA: Modern CUDA Learn Notes with PyTorch for Beginners 🐑</h2>
-    <img src='https://github.com/user-attachments/assets/b2578723-b7a7-4d8f-bcd1-5008947b808a' width="380" height="56" >
+    <img src='https://github.com/user-attachments/assets/b2578723-b7a7-4d8f-bcd1-5008947b808a' width="360" height="56" >
     <a href="https://hellogithub.com/repository/98348655a96640ca8ddcbc298edc901d" target="_blank"><img src="https://api.hellogithub.com/v1/widgets/recommend.svg?rid=98348655a96640ca8ddcbc298edc901d&claim_uid=ofSCbzTmdeQk3FD&theme=dark" alt="Featured｜HelloGitHub" style="width: 250px; height: 54px;" width="250" height="54" /></a>
   </p>
   <div align='center'>

From 0d3e12869026842765462959e6cf40825514923a Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Fri, 27 Jun 2025 19:01:19 +0800
Subject: [PATCH 30/57] fix comments (#345)

---
 kernels/elementwise/elementwise.cu     |  8 +++----
 kernels/elu/elu.cu                     |  8 +++----
 kernels/gelu/gelu.cu                   |  8 +++----
 kernels/hardshrink/hardshrink.cu       | 12 ++++------
 kernels/hardswish/hardswish.cu         | 12 ++++------
 kernels/histogram/histogram.cu         |  3 +--
 kernels/layer-norm/layer_norm.cu       |  8 +++----
 kernels/mat-transpose/mat_transpose.cu | 31 +++++++++-----------------
 kernels/nms/nms.cu                     |  2 --
 kernels/reduce/block_all_reduce.cu     | 20 ++++++++---------
 kernels/relu/relu.cu                   |  9 +++-----
 kernels/rms-norm/rms_norm.cu           |  8 +++----
 kernels/sgemm/sgemm.cu                 |  4 ++--
 kernels/sgemm/sgemm_cublas.cu          |  2 --
 kernels/sgemv/sgemv.cu                 |  6 ++---
 kernels/sigmoid/sigmoid.cu             |  7 +++---
 kernels/softmax/softmax.cu             |  4 ++--
 kernels/swish/swish.cu                 |  9 +++-----
 18 files changed, 64 insertions(+), 97 deletions(-)

diff --git a/kernels/elementwise/elementwise.cu b/kernels/elementwise/elementwise.cu
index 48067371..69e78462 100644
--- a/kernels/elementwise/elementwise.cu
+++ b/kernels/elementwise/elementwise.cu
@@ -17,8 +17,8 @@
 #define BFLOAT2(value) (reinterpret_cast<__nv_bfloat162 *>(&(value))[0])
 #define LDST128BITS(value) (reinterpret_cast<float4 *>(&(value))[0])
 
-// -------------------------------------- FP32
-// -------------------------------------- ElementWise Add grid(N/256),
+// FP32
+// ElementWise Add grid(N/256),
 // block(256) a: Nx1, b: Nx1, c: Nx1, c = elementwise_add(a, b)
 __global__ void elementwise_add_f32_kernel(float *a, float *b, float *c,
                                            int N) {
@@ -45,8 +45,8 @@ __global__ void elementwise_add_f32x4_kernel(float *a, float *b, float *c,
   }
 }
 
-// -------------------------------------- FP16
-// -------------------------------------- ElementWise Add grid(N/256),
+// FP16
+// ElementWise Add grid(N/256),
 // block(256) a: Nx1, b: Nx1, c: Nx1, c = elementwise_add(a, b)
 __global__ void elementwise_add_f16_kernel(half *a, half *b, half *c, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
diff --git a/kernels/elu/elu.cu b/kernels/elu/elu.cu
index 09a819ed..3b71da05 100644
--- a/kernels/elu/elu.cu
+++ b/kernels/elu/elu.cu
@@ -31,13 +31,13 @@
   m.def(STRINGFY(func), &func, STRINGFY(func));
 
 // ELU 计算函数
-// -------------------------------------- FP32
+// FP32
 // --------------------------------------
 __device__ __forceinline__ float elu(float x) {
   return x > 0.f ? x : ALPHA * (expf(x) - 1.f);
 }
 
-// -------------------------------------- FP16
+// FP16
 // --------------------------------------
 __device__ __forceinline__ half elu_half(half x) {
   return __hgt(x, __float2half(0.f))
@@ -46,7 +46,7 @@ __device__ __forceinline__ half elu_half(half x) {
 }
 
 // CUDA 核函数
-// -------------------------------------- FP32
+// FP32
 // --------------------------------------
 __global__ void elu_f32_kernel(float *x, float *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
@@ -67,7 +67,7 @@ __global__ void elu_f32x4_kernel(float *x, float *y, int N) {
   }
 }
 
-// -------------------------------------- FP16
+// FP16
 // --------------------------------------
 __global__ void elu_f16_kernel(half *x, half *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
diff --git a/kernels/gelu/gelu.cu b/kernels/gelu/gelu.cu
index 5cbbf15d..3f7fda9b 100644
--- a/kernels/gelu/gelu.cu
+++ b/kernels/gelu/gelu.cu
@@ -57,8 +57,8 @@ __inline__ __device__ float gelu_none_approximate(float x) {
   return x * 0.5 * (1 + erff(x * M_SQRT1_2));
 }
 
-// -------------------------------------- FP32
-// -------------------------------------- GELU tanh approximate: x, y:x 0.5 * x
+// FP32
+// GELU tanh approximate: x, y:x 0.5 * x
 // * (1.0 + tanh(0.7978845608 * x * (1.0 + 0.044715 * x * x))) grid(N/256),
 // block(K=256)
 __global__ void gelu_f32_kernel(float *x, float *y, int N) {
@@ -91,8 +91,8 @@ __global__ void gelu_f32x4_kernel(float *x, float *y, int N) {
   }
 }
 
-// -------------------------------------- FP16
-// -------------------------------------- GELU approximate: x, y:x 0.5 * x *
+// FP16
+// GELU approximate: x, y:x 0.5 * x *
 // (1.0 + tanh(0.7978845608 (x + 0.044715 * x * x * x))) Vec4
 __global__ void gelu_f16_kernel(half *x, half *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
diff --git a/kernels/hardshrink/hardshrink.cu b/kernels/hardshrink/hardshrink.cu
index 08b0a26d..d391c0d9 100644
--- a/kernels/hardshrink/hardshrink.cu
+++ b/kernels/hardshrink/hardshrink.cu
@@ -31,8 +31,7 @@
   m.def(STRINGFY(func), &func, STRINGFY(func));
 
 // HARDSHRINK 计算函数
-// -------------------------------------- FP32
-// --------------------------------------
+// FP32
 __device__ __forceinline__ float hardshrink(float x) {
   if (x > LAMBD || x < -LAMBD) {
     return x;
@@ -41,8 +40,7 @@ __device__ __forceinline__ float hardshrink(float x) {
   }
 }
 
-// -------------------------------------- FP16
-// --------------------------------------
+// FP16
 __device__ __forceinline__ half hardshrink_half(half x) {
   if (x > __float2half(LAMBD) || x < __float2half(-LAMBD)) {
     return x;
@@ -52,8 +50,7 @@ __device__ __forceinline__ half hardshrink_half(half x) {
 }
 
 // CUDA 核函数
-// -------------------------------------- FP32
-// --------------------------------------
+// FP32
 __global__ void hardshrink_f32_kernel(float *x, float *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   if (idx < N)
@@ -73,8 +70,7 @@ __global__ void hardshrink_f32x4_kernel(float *x, float *y, int N) {
   }
 }
 
-// -------------------------------------- FP16
-// --------------------------------------
+// FP16
 __global__ void hardshrink_f16_kernel(half *x, half *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   if (idx < N)
diff --git a/kernels/hardswish/hardswish.cu b/kernels/hardswish/hardswish.cu
index 35e8769a..7179e57f 100644
--- a/kernels/hardswish/hardswish.cu
+++ b/kernels/hardswish/hardswish.cu
@@ -32,8 +32,7 @@
   m.def(STRINGFY(func), &func, STRINGFY(func));
 
 // HARDSWISH 计算函数
-// -------------------------------------- FP32
-// --------------------------------------
+//  FP32
 __device__ __forceinline__ float hardswish(float x) {
   if (x >= THRESHOLD_A) {
     return x;
@@ -44,8 +43,7 @@ __device__ __forceinline__ float hardswish(float x) {
   }
 }
 
-// -------------------------------------- FP16
-// --------------------------------------
+//  FP16
 __device__ __forceinline__ half hardswish_half(half x) {
   if (x > __float2half(THRESHOLD_A)) {
     return x;
@@ -57,8 +55,7 @@ __device__ __forceinline__ half hardswish_half(half x) {
 }
 
 // CUDA 核函数
-// -------------------------------------- FP32
-// --------------------------------------
+//  FP32
 __global__ void hardswish_f32_kernel(float *x, float *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   if (idx < N)
@@ -78,8 +75,7 @@ __global__ void hardswish_f32x4_kernel(float *x, float *y, int N) {
   }
 }
 
-// -------------------------------------- FP16
-// --------------------------------------
+//  FP16
 __global__ void hardswish_f16_kernel(half *x, half *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   if (idx < N)
diff --git a/kernels/histogram/histogram.cu b/kernels/histogram/histogram.cu
index de61b788..49135a4d 100644
--- a/kernels/histogram/histogram.cu
+++ b/kernels/histogram/histogram.cu
@@ -35,8 +35,7 @@ __global__ void histogram_i32x4_kernel(int *a, int *y, int N) {
   }
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
+// PyTorch bindings for custom kernel
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/layer-norm/layer_norm.cu b/kernels/layer-norm/layer_norm.cu
index 6038d650..cbe2badd 100644
--- a/kernels/layer-norm/layer_norm.cu
+++ b/kernels/layer-norm/layer_norm.cu
@@ -17,8 +17,8 @@
 #define BFLOAT2(value) (reinterpret_cast<__nv_bfloat162 *>(&(value))[0])
 #define LDST128BITS(value) (reinterpret_cast<float4 *>(&(value))[0])
 
-// -------------------------------------- FP32
-// -------------------------------------- Warp Reduce Sum
+//  FP32
+//  Warp Reduce Sum
 template <const int kWarpSize = WARP_SIZE>
 __device__ __forceinline__ float warp_reduce_sum_f32(float val) {
 #pragma unroll
@@ -119,8 +119,8 @@ __global__ void layer_norm_f32x4_kernel(float *x, float *y, float g, float b,
     FLOAT4(y[idx]) = reg_y;
 }
 
-// -------------------------------------- FP16
-// -------------------------------------- Warp Reduce Sum: Half
+//  FP16
+//  Warp Reduce Sum: Half
 template <const int kWarpSize = WARP_SIZE>
 __device__ __forceinline__ half warp_reduce_sum_f16_f16(half val) {
 #pragma unroll
diff --git a/kernels/mat-transpose/mat_transpose.cu b/kernels/mat-transpose/mat_transpose.cu
index 190115c3..d5d9f884 100644
--- a/kernels/mat-transpose/mat_transpose.cu
+++ b/kernels/mat-transpose/mat_transpose.cu
@@ -23,8 +23,8 @@
 #define MAX_EXP_F16 __float2half(11.089866488461016f)
 #define MIN_EXP_F16 __float2half(-9.704060527839234f)
 
-// -------------------------------------- FP32
-// -------------------------------------- col2row means read x[row][col] and
+//  FP32
+//  col2row means read x[row][col] and
 // write y[col][row] row2col means read x[col][row] and write y[row][col]
 __global__ void mat_transpose_f32_col2row_kernel(float *x, float *y,
                                                  const int row, const int col) {
@@ -216,7 +216,6 @@ __global__ void mat_transpose_f32x4_shared_row2col2d_kernel(float *x, float *y,
   }
 }
 
-
 __global__ void mat_transpose_f32x4_shared_bcf_col2row2d_kernel(float *x,
                                                                 float *y,
                                                                 const int row,
@@ -298,11 +297,8 @@ __global__ void mat_transpose_f32x4_shared_bcf_row2col2d_kernel(float *x,
   }
 }
 
-
-__global__ void mat_transpose_f32x4_shared_bcf_merge_write_row2col2d_kernel(float *x,
-                                                                float *y,
-                                                                const int row,
-                                                                const int col) {
+__global__ void mat_transpose_f32x4_shared_bcf_merge_write_row2col2d_kernel(
+    float *x, float *y, const int row, const int col) {
   const int global_x = blockIdx.x * blockDim.x + threadIdx.x;
   const int global_y = blockIdx.y * blockDim.y + threadIdx.y;
   const int local_x = threadIdx.x;
@@ -328,18 +324,13 @@ __global__ void mat_transpose_f32x4_shared_bcf_merge_write_row2col2d_kernel(floa
     smem_val.w = tile[local_x * 4 + 3][local_y];
 
     const int gid_x = blockIdx.x * blockDim.x;
-    const int gid_y = blockIdx.y * blockDim.y * 4;    
+    const int gid_y = blockIdx.y * blockDim.y * 4;
     const int out_y = gid_y + local_x * 4;
     const int out_x = gid_x + local_y;
     reinterpret_cast<float4 *>(y)[(out_x * row + out_y) / 4] = FLOAT4(smem_val);
   }
 }
 
-// TODO: may support double buffer pipeline mat transpose ?
-// TODO: may support fp16 mat transpose ?
-
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
@@ -373,7 +364,7 @@ __global__ void mat_transpose_f32x4_shared_bcf_merge_write_row2col2d_kernel(floa
     dim3 block(WARP_SIZE_S, WARP_SIZE_S);                                      \
     dim3 grid((N + WARP_SIZE_S - 1) / (WARP_SIZE_S * n_element_col),           \
               (M + WARP_SIZE_S - 1) / (WARP_SIZE_S * n_element_row));          \
-    mat_transpose_##tag##2d_kernel <<< grid,                                   \
+    mat_transpose_##tag##2d_kernel < < < grid,                                 \
         block >>> (reinterpret_cast<element_type *>(x.data_ptr()),             \
                    reinterpret_cast<element_type *>(y.data_ptr()), M, N);      \
   }
@@ -400,11 +391,8 @@ TORCH_BINDING_MAT_TRANSPOSE2D(f32x4_shared_bcf_col2row, torch::kFloat32, float,
                               1, 4)
 TORCH_BINDING_MAT_TRANSPOSE2D(f32x4_shared_bcf_row2col, torch::kFloat32, float,
                               4, 1)
-TORCH_BINDING_MAT_TRANSPOSE2D(f32x4_shared_bcf_merge_write_row2col, torch::kFloat32, float,
-                              4, 1)
-
-// TODO: may support double buffer pipeline mat transpose ?
-// TODO: may support fp16 mat transpose ?
+TORCH_BINDING_MAT_TRANSPOSE2D(f32x4_shared_bcf_merge_write_row2col,
+                              torch::kFloat32, float, 4, 1)
 
 // CuTe implentations
 extern void mat_transpose_cute_col2row_reg(torch::Tensor, torch::Tensor);
@@ -442,7 +430,8 @@ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
   // shared memory optimize with bcf
   TORCH_BINDING_COMMON_EXTENSION(mat_transpose_f32x4_shared_bcf_col2row2d)
   TORCH_BINDING_COMMON_EXTENSION(mat_transpose_f32x4_shared_bcf_row2col2d)
-  TORCH_BINDING_COMMON_EXTENSION(mat_transpose_f32x4_shared_bcf_merge_write_row2col2d)
+  TORCH_BINDING_COMMON_EXTENSION(
+      mat_transpose_f32x4_shared_bcf_merge_write_row2col2d)
   // CuTe implentations
   TORCH_BINDING_COMMON_EXTENSION(mat_transpose_cute_col2row_reg)
   TORCH_BINDING_COMMON_EXTENSION(mat_transpose_cute_row2col_reg)
diff --git a/kernels/nms/nms.cu b/kernels/nms/nms.cu
index 95b0b648..f9962c52 100644
--- a/kernels/nms/nms.cu
+++ b/kernels/nms/nms.cu
@@ -58,8 +58,6 @@ __global__ void nms_kernel(const float *boxes, const float *scores, int *keep,
   return;
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/reduce/block_all_reduce.cu b/kernels/reduce/block_all_reduce.cu
index 9f6bb0cc..3932bdf2 100644
--- a/kernels/reduce/block_all_reduce.cu
+++ b/kernels/reduce/block_all_reduce.cu
@@ -25,8 +25,8 @@
 //           MatMul FP8 -> Tensor Cores
 
 // CUDA温故(0x00): 一步步学习block all reduce: 从FP32到FP16/BF16，再到FP8
-// -------------------------------------- FP32
-// -------------------------------------- Warp Reduce Sum
+//  FP32
+//  Warp Reduce Sum
 template <const int kWarpSize = WARP_SIZE>
 __device__ __forceinline__ float warp_reduce_sum_f32(float val) {
 #pragma unroll
@@ -93,8 +93,8 @@ __global__ void block_all_reduce_sum_f32x4_f32_kernel(float *a, float *y,
     atomicAdd(y, sum);
 }
 
-// -------------------------------------- FP16
-// -------------------------------------- Warp Reduce Sum: Half
+//  FP16
+//  Warp Reduce Sum: Half
 template <const int kWarpSize = WARP_SIZE>
 __device__ __forceinline__ half warp_reduce_sum_f16_f16(half val) {
 #pragma unroll
@@ -301,8 +301,8 @@ __global__ void block_all_reduce_sum_f16x8_pack_f32_kernel(half *a, float *y,
     atomicAdd(y, sum);
 }
 
-// -------------------------------------- BF16
-// -------------------------------------- Warp Reduce Sum: Half
+//  BF16
+//  Warp Reduce Sum: Half
 template <const int kWarpSize = WARP_SIZE>
 __device__ __forceinline__ __nv_bfloat16
 warp_reduce_sum_bf16_bf16(__nv_bfloat16 val) {
@@ -520,8 +520,8 @@ __global__ void block_all_reduce_sum_bf16x8_pack_f32_kernel(__nv_bfloat16 *a,
     atomicAdd(y, sum);
 }
 
-// -------------------------------------- FP8
-// --------------------------------------
+//  FP8
+//
 template <const int kWarpSize = WARP_SIZE>
 __device__ __forceinline__ half
 warp_reduce_sum_fp8_e4m3_f16(__nv_fp8_storage_t val) {
@@ -680,8 +680,8 @@ block_all_reduce_sum_fp8_e5m2x16_pack_f16_kernel(__nv_fp8_storage_t *a,
     atomicAdd(y, __half2float(sum));
 }
 
-// -------------------------------------- INT8
-// --------------------------------------
+//  INT8
+//
 template <const int kWarpSize = WARP_SIZE>
 __device__ __forceinline__ int32_t warp_reduce_sum_i8_i32(int8_t val) {
   int32_t val_i32 = static_cast<int32_t>(val);
diff --git a/kernels/relu/relu.cu b/kernels/relu/relu.cu
index 6aaf75a4..ec47ba4d 100644
--- a/kernels/relu/relu.cu
+++ b/kernels/relu/relu.cu
@@ -15,8 +15,8 @@
 #define BFLOAT2(value) (reinterpret_cast<__nv_bfloat162 *>(&(value))[0])
 #define LDST128BITS(value) (reinterpret_cast<float4 *>(&(value))[0])
 
-// -------------------------------------- FP32
-// -------------------------------------- Relu x: N, y: N y=max(0,x)
+//  FP32
+//  Relu x: N, y: N y=max(0,x)
 // grid(N/256), block(K=256)
 __global__ void relu_f32_kernel(float *x, float *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
@@ -39,8 +39,7 @@ __global__ void relu_f32x4_kernel(float *x, float *y, int N) {
   }
 }
 
-// -------------------------------------- FP16
-// --------------------------------------
+//  FP16
 __global__ void relu_f16_kernel(half *x, half *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   if (idx < N)
@@ -106,8 +105,6 @@ __global__ void relu_f16x8_pack_kernel(half *x, half *y, int N) {
   }
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/rms-norm/rms_norm.cu b/kernels/rms-norm/rms_norm.cu
index 1ade3303..1b192e43 100644
--- a/kernels/rms-norm/rms_norm.cu
+++ b/kernels/rms-norm/rms_norm.cu
@@ -17,8 +17,8 @@
 #define BFLOAT2(value) (reinterpret_cast<__nv_bfloat162 *>(&(value))[0])
 #define LDST128BITS(value) (reinterpret_cast<float4 *>(&(value))[0])
 
-// -------------------------------------- FP32
-// -------------------------------------- Warp Reduce Sum
+//  FP32
+//  Warp Reduce Sum
 template <const int kWarpSize = WARP_SIZE>
 __device__ __forceinline__ float warp_reduce_sum_f32(float val) {
 #pragma unroll
@@ -101,8 +101,8 @@ __global__ void rms_norm_f32x4_kernel(float *x, float *y, float g, int N,
     FLOAT4(y[idx]) = reg_y;
 }
 
-// -------------------------------------- FP16
-// -------------------------------------- Warp Reduce Sum: Half
+//  FP16
+//  Warp Reduce Sum: Half
 template <const int kWarpSize = WARP_SIZE>
 __device__ __forceinline__ half warp_reduce_sum_f16_f16(half val) {
 #pragma unroll
diff --git a/kernels/sgemm/sgemm.cu b/kernels/sgemm/sgemm.cu
index c94bd354..9ac15fb8 100644
--- a/kernels/sgemm/sgemm.cu
+++ b/kernels/sgemm/sgemm.cu
@@ -15,8 +15,8 @@
 #define FLOAT4(value) (reinterpret_cast<float4 *>(&(value))[0])
 
 // modified from: https://zhuanlan.zhihu.com/p/657632577
-// -------------------------------------- FP32
-// -------------------------------------- SGEMM naive: compute one c[i,j]
+//  FP32
+//  SGEMM naive: compute one c[i,j]
 // element per threads, all row major
 __global__ void sgemm_naive_f32_kernel(float *a, float *b, float *c, int M,
                                        int N, int K) {
diff --git a/kernels/sgemm/sgemm_cublas.cu b/kernels/sgemm/sgemm_cublas.cu
index 51677a86..7a2d60f7 100644
--- a/kernels/sgemm/sgemm_cublas.cu
+++ b/kernels/sgemm/sgemm_cublas.cu
@@ -41,8 +41,6 @@ void cublas_sgemm_tf32(float *A, float *B, float *C, size_t M, size_t N,
                CUBLAS_GEMM_DEFAULT_TENSOR_OP);
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/sgemv/sgemv.cu b/kernels/sgemv/sgemv.cu
index 0acd2fe8..594a6313 100644
--- a/kernels/sgemv/sgemv.cu
+++ b/kernels/sgemv/sgemv.cu
@@ -14,8 +14,8 @@
 #define INT4(value) (reinterpret_cast<int4 *>(&(value))[0])
 #define FLOAT4(value) (reinterpret_cast<float4 *>(&(value))[0])
 
-// -------------------------------------- FP32
-// -------------------------------------- Warp Reduce Sum
+//  FP32
+//  Warp Reduce Sum
 template <const int kWarpSize = WARP_SIZE>
 __device__ __forceinline__ float warp_reduce_sum_f32(float val) {
 #pragma unroll
@@ -107,8 +107,6 @@ __global__ void sgemv_k16_f32_kernel(float *A, float *x, float *y, int M,
   }
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/sigmoid/sigmoid.cu b/kernels/sigmoid/sigmoid.cu
index c69e6fea..8bab212d 100644
--- a/kernels/sigmoid/sigmoid.cu
+++ b/kernels/sigmoid/sigmoid.cu
@@ -21,8 +21,8 @@
 #define MAX_EXP_F16 __float2half(11.089866488461016f)
 #define MIN_EXP_F16 __float2half(-9.704060527839234f)
 
-// -------------------------------------- FP32
-// -------------------------------------- Sigmoid x: N, y: N y=1/(1+exp(-x))
+//  FP32
+//  Sigmoid x: N, y: N y=1/(1+exp(-x))
 // grid(N/256), block(K=256)
 __global__ void sigmoid_f32_kernel(float *x, float *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
@@ -55,8 +55,7 @@ __global__ void sigmoid_f32x4_kernel(float *x, float *y, int N) {
   }
 }
 
-// -------------------------------------- FP16
-// --------------------------------------
+//  FP16
 __global__ void sigmoid_f16_kernel(half *x, half *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   const half f = __float2half(1.0f);
diff --git a/kernels/softmax/softmax.cu b/kernels/softmax/softmax.cu
index 2279ecd5..35db07a6 100644
--- a/kernels/softmax/softmax.cu
+++ b/kernels/softmax/softmax.cu
@@ -17,8 +17,8 @@
 #define BFLOAT2(value) (reinterpret_cast<__nv_bfloat162 *>(&(value))[0])
 #define LDST128BITS(value) (reinterpret_cast<float4 *>(&(value))[0])
 
-// -------------------------------------- FP32
-// -------------------------------------- DS required for Online Softmax
+//  FP32
+//  DS required for Online Softmax
 struct __align__(8) MD {
   float m;
   float d;
diff --git a/kernels/swish/swish.cu b/kernels/swish/swish.cu
index deefc6f9..d27a9afb 100644
--- a/kernels/swish/swish.cu
+++ b/kernels/swish/swish.cu
@@ -15,8 +15,8 @@
 #define BFLOAT2(value) (reinterpret_cast<__nv_bfloat162 *>(&(value))[0])
 #define LDST128BITS(value) (reinterpret_cast<float4 *>(&(value))[0])
 
-// -------------------------------------- FP32
-// -------------------------------------- Swish x: N, y: N y=x*sigmoid(x)
+//  FP32
+//  Swish x: N, y: N y=x*sigmoid(x)
 __device__ __forceinline__ float swish(float x) {
   return x / (1.0f + expf(-x));
 }
@@ -40,8 +40,7 @@ __global__ void swish_f32x4_kernel(float *x, float *y, int N) {
   }
 }
 
-// -------------------------------------- FP16
-// --------------------------------------
+//  FP16
 __device__ __forceinline__ half swish_half(half x) {
   return __hmul(x, __hdiv(__float2half(1.0f),
                           __hadd(__float2half(1.0f), hexp(__hneg(x)))));
@@ -107,8 +106,6 @@ __global__ void swish_f16x8_pack_kernel(half *x, half *y, int N) {
   }
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));

From edc46f34b4225a9ca1b7ee8c6a1b7ae0e717fe4b Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Fri, 27 Jun 2025 19:18:32 +0800
Subject: [PATCH 31/57] fix comments (#346)

---
 kernels/elementwise/elementwise.cu   | 2 --
 kernels/elu/elu.cu                   | 5 -----
 kernels/embedding/embedding.cu       | 2 --
 kernels/gelu/gelu.cu                 | 2 --
 kernels/hardshrink/hardshrink.cu     | 1 -
 kernels/hardswish/hardswish.cu       | 1 -
 kernels/histogram/histogram.cu       | 1 -
 kernels/layer-norm/layer_norm.cu     | 2 --
 kernels/nvidia-nsight/elementwise.cu | 8 ++++----
 kernels/nvidia-nsight/relu.cu        | 7 +++----
 kernels/reduce/block_all_reduce.cu   | 8 --------
 11 files changed, 7 insertions(+), 32 deletions(-)

diff --git a/kernels/elementwise/elementwise.cu b/kernels/elementwise/elementwise.cu
index 69e78462..1cba70ce 100644
--- a/kernels/elementwise/elementwise.cu
+++ b/kernels/elementwise/elementwise.cu
@@ -120,8 +120,6 @@ __global__ void elementwise_add_f16x8_pack_kernel(half *a, half *b, half *c,
   }
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/elu/elu.cu b/kernels/elu/elu.cu
index 3b71da05..f9cd3c69 100644
--- a/kernels/elu/elu.cu
+++ b/kernels/elu/elu.cu
@@ -32,13 +32,11 @@
 
 // ELU 计算函数
 // FP32
-// --------------------------------------
 __device__ __forceinline__ float elu(float x) {
   return x > 0.f ? x : ALPHA * (expf(x) - 1.f);
 }
 
 // FP16
-// --------------------------------------
 __device__ __forceinline__ half elu_half(half x) {
   return __hgt(x, __float2half(0.f))
              ? x
@@ -47,7 +45,6 @@ __device__ __forceinline__ half elu_half(half x) {
 
 // CUDA 核函数
 // FP32
-// --------------------------------------
 __global__ void elu_f32_kernel(float *x, float *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   if (idx < N)
@@ -68,7 +65,6 @@ __global__ void elu_f32x4_kernel(float *x, float *y, int N) {
 }
 
 // FP16
-// --------------------------------------
 __global__ void elu_f16_kernel(half *x, half *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   if (idx < N)
@@ -129,7 +125,6 @@ __global__ void elu_f16x8_pack_kernel(half *x, half *y, int N) {
   }
 }
 
-// PyTorch 绑定代码
 #define TORCH_BINDING_ELU(packed_type, th_type, element_type, n_elements)      \
   void elu_##packed_type(torch::Tensor x, torch::Tensor y) {                   \
     CHECK_TORCH_TENSOR_DTYPE(x, (th_type))                                     \
diff --git a/kernels/embedding/embedding.cu b/kernels/embedding/embedding.cu
index 074ed264..5b36f757 100644
--- a/kernels/embedding/embedding.cu
+++ b/kernels/embedding/embedding.cu
@@ -78,8 +78,6 @@ __global__ void embedding_f16x8_pack_kernel(const int *idx, half *weight,
       LDST128BITS(weight[offset + 8 * tx]);
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/gelu/gelu.cu b/kernels/gelu/gelu.cu
index 3f7fda9b..822bf887 100644
--- a/kernels/gelu/gelu.cu
+++ b/kernels/gelu/gelu.cu
@@ -182,8 +182,6 @@ __global__ void gelu_f16x8_pack_kernel(half *x, half *y, int N) {
   }
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/hardshrink/hardshrink.cu b/kernels/hardshrink/hardshrink.cu
index d391c0d9..5ea18e8e 100644
--- a/kernels/hardshrink/hardshrink.cu
+++ b/kernels/hardshrink/hardshrink.cu
@@ -131,7 +131,6 @@ __global__ void hardshrink_f16x8_pack_kernel(half *x, half *y, int N) {
   }
 }
 
-// PyTorch 绑定代码
 #define TORCH_BINDING_HARDSHRINK(packed_type, th_type, element_type,           \
                                  n_elements)                                   \
   void hardshrink_##packed_type(torch::Tensor x, torch::Tensor y) {            \
diff --git a/kernels/hardswish/hardswish.cu b/kernels/hardswish/hardswish.cu
index 7179e57f..0c0358a6 100644
--- a/kernels/hardswish/hardswish.cu
+++ b/kernels/hardswish/hardswish.cu
@@ -136,7 +136,6 @@ __global__ void hardswish_f16x8_pack_kernel(half *x, half *y, int N) {
   }
 }
 
-// PyTorch 绑定代码
 #define TORCH_BINDING_HARDSWISH(packed_type, th_type, element_type,            \
                                 n_elements)                                    \
   void hardswish_##packed_type(torch::Tensor x, torch::Tensor y) {             \
diff --git a/kernels/histogram/histogram.cu b/kernels/histogram/histogram.cu
index 49135a4d..334d6207 100644
--- a/kernels/histogram/histogram.cu
+++ b/kernels/histogram/histogram.cu
@@ -35,7 +35,6 @@ __global__ void histogram_i32x4_kernel(int *a, int *y, int N) {
   }
 }
 
-// PyTorch bindings for custom kernel
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/layer-norm/layer_norm.cu b/kernels/layer-norm/layer_norm.cu
index cbe2badd..9962564d 100644
--- a/kernels/layer-norm/layer_norm.cu
+++ b/kernels/layer-norm/layer_norm.cu
@@ -456,8 +456,6 @@ __global__ void layer_norm_f16x8_pack_f32_kernel(half *x, half *y, float g,
   // TODO: support non 8-multiple K here
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/nvidia-nsight/elementwise.cu b/kernels/nvidia-nsight/elementwise.cu
index 8155e946..7322a925 100644
--- a/kernels/nvidia-nsight/elementwise.cu
+++ b/kernels/nvidia-nsight/elementwise.cu
@@ -16,8 +16,8 @@
 #define BFLOAT2(value) (reinterpret_cast<__nv_bfloat162 *>(&(value))[0])
 #define LDST128BITS(value) (reinterpret_cast<float4 *>(&(value))[0])
 
-// -------------------------------------- FP32
-// -------------------------------------- ElementWise Add grid(N/256),
+// FP32
+// ElementWise Add grid(N/256),
 // block(256) a: Nx1, b: Nx1, c: Nx1, c = elementwise_add(a, b)
 __global__ void elementwise_add_f32_kernel(float *a, float *b, float *c,
                                            int N) {
@@ -44,8 +44,8 @@ __global__ void elementwise_add_f32x4_kernel(float *a, float *b, float *c,
   }
 }
 
-// -------------------------------------- FP16
-// -------------------------------------- ElementWise Add grid(N/256),
+// FP16
+// ElementWise Add grid(N/256),
 // block(256) a: Nx1, b: Nx1, c: Nx1, c = elementwise_add(a, b)
 __global__ void elementwise_add_f16_kernel(half *a, half *b, half *c, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
diff --git a/kernels/nvidia-nsight/relu.cu b/kernels/nvidia-nsight/relu.cu
index bed7054b..9a830f75 100644
--- a/kernels/nvidia-nsight/relu.cu
+++ b/kernels/nvidia-nsight/relu.cu
@@ -14,8 +14,8 @@
 #define BFLOAT2(value) (reinterpret_cast<__nv_bfloat162 *>(&(value))[0])
 #define LDST128BITS(value) (reinterpret_cast<float4 *>(&(value))[0])
 
-// -------------------------------------- FP32
-// -------------------------------------- Relu x: N, y: N y=max(0,x)
+// FP32
+// Relu x: N, y: N y=max(0,x)
 // grid(N/256), block(K=256)
 __global__ void relu_f32_kernel(float *x, float *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
@@ -38,8 +38,7 @@ __global__ void relu_f32x4_kernel(float *x, float *y, int N) {
   }
 }
 
-// -------------------------------------- FP16
-// --------------------------------------
+// FP16
 __global__ void relu_f16_kernel(half *x, half *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   if (idx < N)
diff --git a/kernels/reduce/block_all_reduce.cu b/kernels/reduce/block_all_reduce.cu
index 3932bdf2..e6000ae7 100644
--- a/kernels/reduce/block_all_reduce.cu
+++ b/kernels/reduce/block_all_reduce.cu
@@ -17,14 +17,6 @@
 #define BFLOAT2(value) (reinterpret_cast<__nv_bfloat162 *>(&(value))[0])
 #define LDST128BITS(value) (reinterpret_cast<float4 *>(&(value))[0])
 
-// FP16/BF16 CUDA Cores/Tensor Cores:
-// https://resources.nvidia.com/en-us-tensor-core
-// Non MatMul FP16/BF16 -> CUDA Cores
-//     MatMul FP16/BF16 -> Tensor Cores
-//       Non MatMul FP8 -> Not supported
-//           MatMul FP8 -> Tensor Cores
-
-// CUDA温故(0x00): 一步步学习block all reduce: 从FP32到FP16/BF16，再到FP8
 //  FP32
 //  Warp Reduce Sum
 template <const int kWarpSize = WARP_SIZE>

From 04e331fa254aa7ae9440ff34bbc30268e26dd1df Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Fri, 27 Jun 2025 21:59:14 +0800
Subject: [PATCH 32/57] fix rope.cu

---
 kernels/rope/rope.cu | 2 --
 1 file changed, 2 deletions(-)

diff --git a/kernels/rope/rope.cu b/kernels/rope/rope.cu
index 63422f99..d4367b72 100644
--- a/kernels/rope/rope.cu
+++ b/kernels/rope/rope.cu
@@ -67,8 +67,6 @@ __global__ void rope_f32x4_pack_kernel(float *x, float *out, int seq_len,
   FLOAT4(out[idx * 4]) = out_v;
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));

From 8c5464d26c8d12cc4d78e2a52924957644a53d4e Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Fri, 27 Jun 2025 22:04:50 +0800
Subject: [PATCH 33/57] fix all2all example (#347)

---
 others/pytorch/distributed/test_all_to_all_single_ray.py | 3 ---
 1 file changed, 3 deletions(-)

diff --git a/others/pytorch/distributed/test_all_to_all_single_ray.py b/others/pytorch/distributed/test_all_to_all_single_ray.py
index 9169b9f8..d0334364 100644
--- a/others/pytorch/distributed/test_all_to_all_single_ray.py
+++ b/others/pytorch/distributed/test_all_to_all_single_ray.py
@@ -129,7 +129,4 @@ def run(self) -> torch.Tensor:
     ]
     results = ray.get([worker.run.remote() for worker in workers])
     ray.shutdown()
-    if dist.is_initialized():
-        dist.destroy_process_group()
-        print("Process group destroyed")
     print("All processes completed successfully")

From d9eb9a58e1b792f462e75bee7c0169ac10747c5c Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Fri, 27 Jun 2025 22:06:33 +0800
Subject: [PATCH 34/57] fix softmax.cu (#348)

---
 kernels/softmax/softmax.cu | 10 ++++------
 1 file changed, 4 insertions(+), 6 deletions(-)

diff --git a/kernels/softmax/softmax.cu b/kernels/softmax/softmax.cu
index 35db07a6..d5c7ed9f 100644
--- a/kernels/softmax/softmax.cu
+++ b/kernels/softmax/softmax.cu
@@ -17,8 +17,8 @@
 #define BFLOAT2(value) (reinterpret_cast<__nv_bfloat162 *>(&(value))[0])
 #define LDST128BITS(value) (reinterpret_cast<float4 *>(&(value))[0])
 
-//  FP32
-//  DS required for Online Softmax
+// FP32
+// DS required for Online Softmax
 struct __align__(8) MD {
   float m;
   float d;
@@ -102,8 +102,8 @@ __device__ float block_reduce_max_f32(float val) {
   return value;
 }
 
-// // Softmax x: N, y: N
-// // grid(N/256), block(K=256)
+// Softmax x: N, y: N
+// grid(N/256), block(K=256)
 // template<const int NUM_THREADS = 256>
 // __global__ void softmax_f32_kernel(float* x, float* y, float* total, int N) {
 
@@ -408,8 +408,6 @@ online_safe_softmax_f32x4_pack_per_token_kernel(float *x, float *y, int N) {
   }
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));

From 0ba539003a983b75e218c10d2ccd70b03d73cdfb Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Fri, 27 Jun 2025 22:07:42 +0800
Subject: [PATCH 35/57] fix CONTRIBUTE.md (#349)

---
 CONTRIBUTE.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/CONTRIBUTE.md b/CONTRIBUTE.md
index c2ba5e5c..f1bbfa12 100644
--- a/CONTRIBUTE.md
+++ b/CONTRIBUTE.md
@@ -1,7 +1,7 @@
 # Developer Guide
 
 ## 🌤🌤 Goals
-Any kernel implementation is welcome. This repository is primarily for learning/practice, and achieving optimal performance is not the final goal here—focus on "using first, then using well". For optimal performance, it is recommended to directly use official implementations such as cuBLAS, cuDNN, FlashAttention, TensorRT, etc. If you are interested in implementing a specific kernel in this repository, feel free to send a new PR to LeetCUDA, for example:
+Any kernel implementation is welcome. This repository is primarily for learning/practice, and achieving optimal performance is not the final goal. LeetCUDA focus on "using first, then using well". For optimal performance, it is recommended to directly use official implementations such as cuBLAS, cuDNN, FlashAttention, TensorRT, etc. If you are interested in implementing a specific kernel in this repository, feel free to send a new PR to LeetCUDA, for example:
 
 ## 👨‍💻👨‍💻 Pre-commit
 

From 497b8c352bb42347f39a450cffdfebafb3dd84a8 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Sat, 28 Jun 2025 12:25:28 +0800
Subject: [PATCH 36/57] feat: add torch.compile blogs (#350)

* fix comments

* fix comments

* fix comments

* fix comments

* add torch.compile blogs

* add torch.compile blogs

* add torch.compile blogs
---
 .pre-commit-config.yaml                       |  2 -
 README.md                                     | 54 +++++--------------
 kernels/dot-product/dot_product.cu            | 10 ++--
 kernels/hardswish/hardswish.cu                | 12 ++---
 kernels/hgemm/cublas/hgemm_cublas.cu          |  3 +-
 .../hgemm/cutlass/hgemm_mma_stage_tn_cute.cu  |  3 +-
 kernels/hgemm/mma/basic/hgemm_mma.cu          |  2 -
 kernels/hgemm/mma/basic/hgemm_mma_stage.cu    |  2 -
 kernels/hgemm/mma/basic/hgemm_mma_stage_tn.cu |  2 -
 .../mma/swizzle/hgemm_mma_stage_swizzle.cu    |  2 -
 .../mma/swizzle/hgemm_mma_stage_tn_swizzle.cu |  3 +-
 .../swizzle/hgemm_mma_stage_tn_swizzle_x2.cu  |  3 +-
 .../swizzle/hgemm_mma_stage_tn_swizzle_x4.cu  |  3 +-
 kernels/hgemm/naive/hgemm.cu                  |  6 +--
 kernels/hgemm/naive/hgemm_async.cu            |  2 -
 kernels/hgemv/hgemv.cu                        |  6 +--
 kernels/hgemv/hgemv_cute.cu                   |  2 -
 kernels/layer-norm/layer_norm.cu              |  8 +--
 kernels/mat-transpose/mat_transpose.cu        |  4 +-
 kernels/reduce/block_all_reduce.cu            |  6 +--
 kernels/relu/relu.cu                          |  4 +-
 kernels/rms-norm/rms_norm.cu                  |  2 -
 kernels/sgemm/sgemm.cu                        |  6 +--
 kernels/sgemm/sgemm_async.cu                  |  2 -
 kernels/sgemm/sgemm_wmma_tf32_stage.cu        |  2 -
 kernels/sigmoid/sigmoid.cu                    |  6 +--
 kernels/swish/swish.cu                        |  4 +-
 27 files changed, 46 insertions(+), 115 deletions(-)

diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml
index d4be4b47..cb5d080d 100644
--- a/.pre-commit-config.yaml
+++ b/.pre-commit-config.yaml
@@ -10,8 +10,6 @@ repos:
       - id: end-of-file-fixer
       - id: check-yaml
         args: [--allow-multiple-documents]
-      - id: check-toml
-      - id: check-ast
       - id: check-added-large-files
       - id: check-merge-conflict
       - id: check-shebang-scripts-are-executable
diff --git a/README.md b/README.md
index 6d307311..d97af656 100644
--- a/README.md
+++ b/README.md
@@ -1,15 +1,3 @@
-<!---
-  <img src='https://github.com/user-attachments/assets/9306862b-2a30-4a87-bb33-0fde9e9d7cea' width=250 >
-      <a href="#cuda-kernel">📚200+ CUDA Kernels</a> | <a href="#my-blogs-part-1"> 📚100+ LLM/CUDA Blogs</a> | <a href="#HGEMM-bench"> ⚡️HGEMM MMA</a> | <a href="#fa-mma-bench"> ⚡️FA-2 MMA </a> <p>
-<img src='https://github.com/user-attachments/assets/b2578723-b7a7-4d8f-bcd1-5008947b808a' >
-<div align="center">
-  <p align="center">
-    <a href="#contribute">愿以青衿涉险苦，为君先踏棘荆途。他年若览通衢阔，莫忘初逢问路吾。</a>
-  </p>
-</div>
---->
-
-
 <div align="center">
   <p align="center">
     <h2>📚 LeetCUDA: Modern CUDA Learn Notes with PyTorch for Beginners 🐑</h2>
@@ -24,14 +12,14 @@
       <img src=https://img.shields.io/github/stars/xlite-dev/LeetCUDA.svg?style=social >
       <img src=https://img.shields.io/badge/Release-v3.0.6-brightgreen.svg >
       <img src=https://img.shields.io/badge/License-GPLv3.0-turquoise.svg >
- </div>
+  </div>
 </div>
 
 📚 **LeetCUDA**: It includes **Tensor/CUDA Cores, TF32/F16/BF16/F8**, [📖200+ CUDA Kernels🔥](#cuda-kernel) with PyTorch, [📖100+ LLM/CUDA🔥](#my-blogs-part-1) blogs, [📖HGEMM⚡️](./kernels/hgemm) which can achieve `98%~100%` TFLOPS of **cuBLAS**, and [📖flash-attn⚡️](./kernels/flash-attn) using Tensor Cores with pure MMA PTX. ♥️ Please consider to leave a ⭐️ Star to support me, my bro ~ ♥️
 
 <div align="center">
   <p align="center">
-    <a href="#contribute">🔥🔥 PR Welcome: Add Your Kernel to LeetCUDA! Let's make it Awesome together! 🎉🎉</a>
+    <a href="#contribute">🔥🔥 PR Welcome: Add Your Kernel to LeetCUDA! Let's make it Awesome together! 🎉🎉</a> <br>
     <a href=https://github.com/xlite-dev/LeetCUDA/graphs/contributors > <img src=https://opencollective.com/leetcuda/contributors.svg height=40px > </a>
   </p>
 </div>
@@ -52,7 +40,7 @@
 ## 📖 News 🔥🔥
 <div id="news"></div>
 
-- [2025-06-16]: [🤗CacheDiT](https://github.com/vipshop/cache-dit) is release! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for Diffusion Transformers (**DBCache**, **DBPrune**, **FBCache**, etc)🔥. Feel free to take a try!  
+- [2025-06-16]: [🤗CacheDiT](https://github.com/vipshop/cache-dit) is release! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for Diffusion Transformers (**DBCache**, **DBPrune**, **FBCache**, etc)🔥. Feel free to take a try!
 
 <div align='center'>
   <img src='https://github.com/user-attachments/assets/a5ec4320-d2f9-4254-888a-170b2d9e3784' height=170px>
@@ -77,31 +65,6 @@
 
 ## 📖 Contents
 <div id="contents"></div>
-<!---
-- [📖 HGEMM-MMA 🎉🎉](#HGEMM-bench)
-  - [📚 CUDA/Tensor Cores](#HGEMM-bench)
-  - [📚 Tile Block(Br, Bc)](#HGEMM-bench)
-  - [📚 Tile MMAs/Warps](#HGEMM-bench)
-  - [📚 Pack LDST(128 bits)](#HGEMM-bench)
-  - [📚 Multi Stages(2~4)](#HGEMM-bench)
-  - [📚 Block/Warp Swizzle](#HGEMM-bench)
-  - [📚 SMEM Swizzle](#HGEMM-bench)
-  - [📚 Register Double Buffers](#HGEMM-bench)
-  - [📚 Collective Store(Shfl)](#HGEMM-bench)
-  - [📚 Layout NN/TN](#HGEMM-bench)
-- [📖 FlashAttention-MMA 🎉🎉](#fa-mma-bench)
-- [📖 200+ CUDA Kernels 🔥🔥](#cuda-kernel)
-- [📖 100+ 高性能计算文章 💡💡](#my-blogs-part-1)
-  - [📚 大模型推理优化原理](#my-blogs-part-1)
-  - [📚 大模型分布式训推原理](#my-blogs-part-1)
-  - [📚 CV/C++/模型部署优化](#my-blogs-part-1)
-  - [📚 CUDA优化入门与实践](#other-blogs)
-  - [📚 Tensor Cores入门教程](#other-blogs)
-  - [📚 CuTe系列详解与实践](#other-blogs)
-  - [📚 GPU指令集架构精解](#other-blogs)
-  - [📚 GPU通信架构精解](#other-blogs)
-- [📖 How to Contribute 👀👇](#contribute)
---->
 
 - [📖 HGEMM-MMA 🎉🎉](#HGEMM-bench)
 - [📖 FlashAttention-MMA 🎉🎉](#fa-mma-bench)
@@ -521,7 +484,7 @@ The kernels listed here will guide you through a step-by-step progression, rangi
 
 |📖 类型-标题|📖 作者| 📖 推荐 |
 |:---|:---|:---|
-| [[Diffusion推理]📖DiT推理加速综述: Caching](https://zhuanlan.zhihu.com/p/711223667)|@DefTruth|⭐️⭐️⭐|  
+| [[Diffusion推理]📖DiT推理加速综述: Caching](https://zhuanlan.zhihu.com/p/711223667)|@DefTruth|⭐️⭐️⭐|
 | [[Triton编程][基础]📖Triton极简入门: Triton Vector Add](https://zhuanlan.zhihu.com/p/1902778199261291694)|@DefTruth|⭐️⭐️⭐|
 | [[Triton编程][基础]📖Triton Fused Softmax Kernel详解: 从Python源码到PTX](https://zhuanlan.zhihu.com/p/1899562146477609112)|@DefTruth|⭐️⭐️⭐|
 | [[Triton编程][基础]📖vLLM Triton Merge Attention States Kernel详解](https://zhuanlan.zhihu.com/p/1904937907703243110)|@DefTruth|⭐️⭐️⭐|
@@ -665,6 +628,15 @@ The kernels listed here will guide you through a step-by-step progression, rangi
 | [[Tensor Cores]📖Nvidia Tensor Core-MMA PTX编程入门](https://zhuanlan.zhihu.com/p/621855199)|@木子知|⭐️⭐️⭐️|
 | [[Tensor Cores]📖CUDA Ampere Tensor Core HGEMM 矩阵乘法优化](https://zhuanlan.zhihu.com/p/555339335)|@nicholaswilde|⭐️⭐️⭐️|
 | [[GPU通信架构][精解]📖NVIDIA GPGPU（四）- 通信架构](https://zhuanlan.zhihu.com/p/680262016)|@Bruce|⭐️⭐️⭐️|
+| [[torch.compile][原理]📖Torch.compile流程解析: 介绍](https://zhuanlan.zhihu.com/p/9418379234)|@StarCap|⭐️⭐️⭐️|
+| [[torch.compile][原理]📖Torch.compile流程解析: TorchDynamo](https://zhuanlan.zhihu.com/p/9640728231)|@StarCap|⭐️⭐️⭐️|
+| [[torch.compile][原理]📖Torch.compile流程解析: AOTAutograd](https://zhuanlan.zhihu.com/p/9997263922)|@StarCap|⭐️⭐️⭐️|
+| [[torch.compile][原理]📖Torch.compile流程解析: TorchInductor](https://zhuanlan.zhihu.com/p/11224299472)|@StarCap|⭐️⭐️⭐️|
+| [[torch.compile][原理]📖Torch.compile流程解析: 算子融合](https://zhuanlan.zhihu.com/p/21053905491)|@StarCap|⭐️⭐️⭐️|
+| [[torch.compile][实践]📖Torch.compile使用指南](https://zhuanlan.zhihu.com/p/620163218)|@jhang|⭐️⭐️⭐️|
+| [[torch.compile][实践]📖Torch.compile详细示例解析教程](https://zhuanlan.zhihu.com/p/855291863)|@Bbuf|⭐️⭐️⭐️|
+| [[torch.compile][原理]📖一文搞懂TorchDynamo原理](https://zhuanlan.zhihu.com/p/630933479)|@吾乃阿尔法|⭐️⭐️⭐️|
+| [[torch.compile][原理]📖理解torch.compile基本原理和使用方式](https://zhuanlan.zhihu.com/p/12712224407)|@俯仰|⭐️⭐️⭐️|
 
 ## ©️License ([©️back👆🏻](#contents))
 
diff --git a/kernels/dot-product/dot_product.cu b/kernels/dot-product/dot_product.cu
index c921505b..12ea4195 100644
--- a/kernels/dot-product/dot_product.cu
+++ b/kernels/dot-product/dot_product.cu
@@ -17,8 +17,8 @@
 #define BFLOAT2(value) (reinterpret_cast<__nv_bfloat162 *>(&(value))[0])
 #define LDST128BITS(value) (reinterpret_cast<float4 *>(&(value))[0])
 
-// -------------------------------------- FP32
-// -------------------------------------- Warp Reduce Sum
+// FP32
+// Warp Reduce Sum
 template <const int kWarpSize = WARP_SIZE>
 __device__ __forceinline__ float warp_reduce_sum_f32(float val) {
 #pragma unroll
@@ -87,8 +87,8 @@ __global__ void dot_prod_f32x4_f32_kernel(float *a, float *b, float *y, int N) {
     atomicAdd(y, prod);
 }
 
-// -------------------------------------- FP16
-// -------------------------------------- Warp Reduce Sum: Half
+// FP16
+// Warp Reduce Sum: Half
 template <const int kWarpSize = WARP_SIZE>
 __device__ __forceinline__ half warp_reduce_sum_f16_f16(half val) {
 #pragma unroll
@@ -199,8 +199,6 @@ __global__ void dot_prod_f16x8_pack_f32_kernel(half *a, half *b, float *y,
     atomicAdd(y, prod);
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/hardswish/hardswish.cu b/kernels/hardswish/hardswish.cu
index 0c0358a6..8c0f641e 100644
--- a/kernels/hardswish/hardswish.cu
+++ b/kernels/hardswish/hardswish.cu
@@ -19,20 +19,17 @@
 #define BFLOAT2(value) (reinterpret_cast<__nv_bfloat162 *>(&(value))[0])
 #define LDST128BITS(value) (reinterpret_cast<float4 *>(&(value))[0])
 
-// 定义 CHECK_TORCH_TENSOR_DTYPE 宏
 #define CHECK_TORCH_TENSOR_DTYPE(T, th_type)                                   \
   if (((T).options().dtype() != (th_type))) {                                  \
     std::cout << "Tensor Info:" << (T).options() << std::endl;                 \
     throw std::runtime_error("Tensor dtype must be " #th_type);                \
   }
 
-// 定义 TORCH_BINDING_COMMON_EXTENSION 宏
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
 
-// HARDSWISH 计算函数
-//  FP32
+// FP32
 __device__ __forceinline__ float hardswish(float x) {
   if (x >= THRESHOLD_A) {
     return x;
@@ -43,7 +40,7 @@ __device__ __forceinline__ float hardswish(float x) {
   }
 }
 
-//  FP16
+// FP16
 __device__ __forceinline__ half hardswish_half(half x) {
   if (x > __float2half(THRESHOLD_A)) {
     return x;
@@ -54,8 +51,7 @@ __device__ __forceinline__ half hardswish_half(half x) {
   }
 }
 
-// CUDA 核函数
-//  FP32
+// FP32
 __global__ void hardswish_f32_kernel(float *x, float *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   if (idx < N)
@@ -75,7 +71,7 @@ __global__ void hardswish_f32x4_kernel(float *x, float *y, int N) {
   }
 }
 
-//  FP16
+// FP16
 __global__ void hardswish_f16_kernel(half *x, half *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   if (idx < N)
diff --git a/kernels/hgemm/cublas/hgemm_cublas.cu b/kernels/hgemm/cublas/hgemm_cublas.cu
index f7d63be2..356319f5 100644
--- a/kernels/hgemm/cublas/hgemm_cublas.cu
+++ b/kernels/hgemm/cublas/hgemm_cublas.cu
@@ -173,8 +173,7 @@ int main(int argc, char *argv[]) {
 }
 // build torch python binding
 #else
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
+
 #include <torch/extension.h>
 #include <torch/types.h>
 
diff --git a/kernels/hgemm/cutlass/hgemm_mma_stage_tn_cute.cu b/kernels/hgemm/cutlass/hgemm_mma_stage_tn_cute.cu
index b3e6354d..b6757a20 100644
--- a/kernels/hgemm/cutlass/hgemm_mma_stage_tn_cute.cu
+++ b/kernels/hgemm/cutlass/hgemm_mma_stage_tn_cute.cu
@@ -469,8 +469,7 @@ int main() {
 
 #include <torch/extension.h>
 #include <torch/types.h>
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
+
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/hgemm/mma/basic/hgemm_mma.cu b/kernels/hgemm/mma/basic/hgemm_mma.cu
index af2c83b3..6e34c036 100644
--- a/kernels/hgemm/mma/basic/hgemm_mma.cu
+++ b/kernels/hgemm/mma/basic/hgemm_mma.cu
@@ -288,8 +288,6 @@ __global__ void __launch_bounds__(256)
   }
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/hgemm/mma/basic/hgemm_mma_stage.cu b/kernels/hgemm/mma/basic/hgemm_mma_stage.cu
index b8117969..22669ea1 100644
--- a/kernels/hgemm/mma/basic/hgemm_mma_stage.cu
+++ b/kernels/hgemm/mma/basic/hgemm_mma_stage.cu
@@ -2039,8 +2039,6 @@ int main(int argc, char *argv[]) {
 
 #else
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #include <torch/extension.h>
 #include <torch/types.h>
 #define STRINGFY(str) #str
diff --git a/kernels/hgemm/mma/basic/hgemm_mma_stage_tn.cu b/kernels/hgemm/mma/basic/hgemm_mma_stage_tn.cu
index 27f70b9a..0a7ecdc2 100644
--- a/kernels/hgemm/mma/basic/hgemm_mma_stage_tn.cu
+++ b/kernels/hgemm/mma/basic/hgemm_mma_stage_tn.cu
@@ -492,8 +492,6 @@ int main(int argc, char *argv[]) {
 
 #else
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #include <torch/extension.h>
 #include <torch/types.h>
 #define STRINGFY(str) #str
diff --git a/kernels/hgemm/mma/swizzle/hgemm_mma_stage_swizzle.cu b/kernels/hgemm/mma/swizzle/hgemm_mma_stage_swizzle.cu
index e098d288..943e869b 100644
--- a/kernels/hgemm/mma/swizzle/hgemm_mma_stage_swizzle.cu
+++ b/kernels/hgemm/mma/swizzle/hgemm_mma_stage_swizzle.cu
@@ -738,8 +738,6 @@ int main(int argc, char *argv[]) {
 
 #else
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #include <torch/extension.h>
 #include <torch/types.h>
 #define STRINGFY(str) #str
diff --git a/kernels/hgemm/mma/swizzle/hgemm_mma_stage_tn_swizzle.cu b/kernels/hgemm/mma/swizzle/hgemm_mma_stage_tn_swizzle.cu
index 526048ba..a45b41a2 100644
--- a/kernels/hgemm/mma/swizzle/hgemm_mma_stage_tn_swizzle.cu
+++ b/kernels/hgemm/mma/swizzle/hgemm_mma_stage_tn_swizzle.cu
@@ -569,8 +569,7 @@ int main(int argc, char *argv[]) {
 }
 
 #else
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
+
 #include <torch/extension.h>
 #include <torch/types.h>
 #define STRINGFY(str) #str
diff --git a/kernels/hgemm/mma/swizzle/hgemm_mma_stage_tn_swizzle_x2.cu b/kernels/hgemm/mma/swizzle/hgemm_mma_stage_tn_swizzle_x2.cu
index 4d2aca80..cbe65a40 100644
--- a/kernels/hgemm/mma/swizzle/hgemm_mma_stage_tn_swizzle_x2.cu
+++ b/kernels/hgemm/mma/swizzle/hgemm_mma_stage_tn_swizzle_x2.cu
@@ -752,8 +752,7 @@ int main(int argc, char *argv[]) {
 }
 
 #else
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
+
 #include <torch/extension.h>
 #include <torch/types.h>
 #define STRINGFY(str) #str
diff --git a/kernels/hgemm/mma/swizzle/hgemm_mma_stage_tn_swizzle_x4.cu b/kernels/hgemm/mma/swizzle/hgemm_mma_stage_tn_swizzle_x4.cu
index 6ec8d7e1..2b19e9e6 100644
--- a/kernels/hgemm/mma/swizzle/hgemm_mma_stage_tn_swizzle_x4.cu
+++ b/kernels/hgemm/mma/swizzle/hgemm_mma_stage_tn_swizzle_x4.cu
@@ -823,8 +823,7 @@ int main(int argc, char *argv[]) {
 }
 
 #else
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
+
 #include <torch/extension.h>
 #include <torch/types.h>
 #define STRINGFY(str) #str
diff --git a/kernels/hgemm/naive/hgemm.cu b/kernels/hgemm/naive/hgemm.cu
index a0593c02..50d07377 100644
--- a/kernels/hgemm/naive/hgemm.cu
+++ b/kernels/hgemm/naive/hgemm.cu
@@ -18,8 +18,8 @@
 #define LDST64BITS(value) (reinterpret_cast<float2 *>(&(value))[0])
 #define LDST128BITS(value) (reinterpret_cast<float4 *>(&(value))[0])
 
-// -------------------------------------- FP16
-// -------------------------------------- HGEMM naive: compute one c[i,j]
+// FP16
+// HGEMM naive: compute one c[i,j]
 // element per threads, all row major
 __global__ void hgemm_naive_f16_kernel(half *a, half *b, half *c, int M, int N,
                                        int K) {
@@ -789,8 +789,6 @@ __global__ void hgemm_t_8x8_sliced_k_f16x8_pack_bcf_dbuf_kernel(
   }
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/hgemm/naive/hgemm_async.cu b/kernels/hgemm/naive/hgemm_async.cu
index 4ff65c41..139712ec 100644
--- a/kernels/hgemm/naive/hgemm_async.cu
+++ b/kernels/hgemm/naive/hgemm_async.cu
@@ -743,8 +743,6 @@ hgemm_t_16x8_sliced_k32_f16x8_pack_dbuf_async_kernel(half *a, half *b, half *c,
   }
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/hgemv/hgemv.cu b/kernels/hgemv/hgemv.cu
index 08b7e1cd..3f2f7f01 100644
--- a/kernels/hgemv/hgemv.cu
+++ b/kernels/hgemv/hgemv.cu
@@ -16,8 +16,8 @@
 #define HALF2(value) (reinterpret_cast<half2 *>(&(value))[0])
 #define BFLOAT2(value) (reinterpret_cast<__nv_bfloat162 *>(&(value))[0])
 
-// -------------------------------------- FP16
-// -------------------------------------- Warp Reduce Sum
+// FP16
+// Warp Reduce Sum
 template <const int kWarpSize = WARP_SIZE>
 __device__ __forceinline__ half warp_reduce_sum_f16(half val) {
 #pragma unroll
@@ -109,8 +109,6 @@ __global__ void hgemv_k16_f16_kernel(half *A, half *x, half *y, int M, int K) {
   }
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/hgemv/hgemv_cute.cu b/kernels/hgemv/hgemv_cute.cu
index 8ad9761d..f1691130 100644
--- a/kernels/hgemv/hgemv_cute.cu
+++ b/kernels/hgemv/hgemv_cute.cu
@@ -284,8 +284,6 @@ __global__ void hgemv_tensor_core_cute_kernel(typename HgemvConfig_::T *Aptr,
   copy_if(elem_pred2, sum, gC(elem_index2, _));
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/layer-norm/layer_norm.cu b/kernels/layer-norm/layer_norm.cu
index 9962564d..20a43600 100644
--- a/kernels/layer-norm/layer_norm.cu
+++ b/kernels/layer-norm/layer_norm.cu
@@ -17,8 +17,8 @@
 #define BFLOAT2(value) (reinterpret_cast<__nv_bfloat162 *>(&(value))[0])
 #define LDST128BITS(value) (reinterpret_cast<float4 *>(&(value))[0])
 
-//  FP32
-//  Warp Reduce Sum
+// FP32
+// Warp Reduce Sum
 template <const int kWarpSize = WARP_SIZE>
 __device__ __forceinline__ float warp_reduce_sum_f32(float val) {
 #pragma unroll
@@ -119,8 +119,8 @@ __global__ void layer_norm_f32x4_kernel(float *x, float *y, float g, float b,
     FLOAT4(y[idx]) = reg_y;
 }
 
-//  FP16
-//  Warp Reduce Sum: Half
+// FP16
+// Warp Reduce Sum: Half
 template <const int kWarpSize = WARP_SIZE>
 __device__ __forceinline__ half warp_reduce_sum_f16_f16(half val) {
 #pragma unroll
diff --git a/kernels/mat-transpose/mat_transpose.cu b/kernels/mat-transpose/mat_transpose.cu
index d5d9f884..d3dd6eff 100644
--- a/kernels/mat-transpose/mat_transpose.cu
+++ b/kernels/mat-transpose/mat_transpose.cu
@@ -23,8 +23,8 @@
 #define MAX_EXP_F16 __float2half(11.089866488461016f)
 #define MIN_EXP_F16 __float2half(-9.704060527839234f)
 
-//  FP32
-//  col2row means read x[row][col] and
+// FP32
+// col2row means read x[row][col] and
 // write y[col][row] row2col means read x[col][row] and write y[row][col]
 __global__ void mat_transpose_f32_col2row_kernel(float *x, float *y,
                                                  const int row, const int col) {
diff --git a/kernels/reduce/block_all_reduce.cu b/kernels/reduce/block_all_reduce.cu
index e6000ae7..5d56e464 100644
--- a/kernels/reduce/block_all_reduce.cu
+++ b/kernels/reduce/block_all_reduce.cu
@@ -17,8 +17,8 @@
 #define BFLOAT2(value) (reinterpret_cast<__nv_bfloat162 *>(&(value))[0])
 #define LDST128BITS(value) (reinterpret_cast<float4 *>(&(value))[0])
 
-//  FP32
-//  Warp Reduce Sum
+// FP32
+// Warp Reduce Sum
 template <const int kWarpSize = WARP_SIZE>
 __device__ __forceinline__ float warp_reduce_sum_f32(float val) {
 #pragma unroll
@@ -751,8 +751,6 @@ __global__ void block_all_reduce_sum_i8x16_pack_i32_kernel(int8_t *a,
     atomicAdd(y, sum);
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/relu/relu.cu b/kernels/relu/relu.cu
index ec47ba4d..4569c569 100644
--- a/kernels/relu/relu.cu
+++ b/kernels/relu/relu.cu
@@ -15,8 +15,8 @@
 #define BFLOAT2(value) (reinterpret_cast<__nv_bfloat162 *>(&(value))[0])
 #define LDST128BITS(value) (reinterpret_cast<float4 *>(&(value))[0])
 
-//  FP32
-//  Relu x: N, y: N y=max(0,x)
+// FP32
+// Relu x: N, y: N y=max(0,x)
 // grid(N/256), block(K=256)
 __global__ void relu_f32_kernel(float *x, float *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
diff --git a/kernels/rms-norm/rms_norm.cu b/kernels/rms-norm/rms_norm.cu
index 1b192e43..aeee45eb 100644
--- a/kernels/rms-norm/rms_norm.cu
+++ b/kernels/rms-norm/rms_norm.cu
@@ -413,8 +413,6 @@ __global__ void rms_norm_f16x8_pack_f32_kernel(half *x, half *y, float g, int N,
   // TODO: support non 8-multiple K here
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/sgemm/sgemm.cu b/kernels/sgemm/sgemm.cu
index 9ac15fb8..50997504 100644
--- a/kernels/sgemm/sgemm.cu
+++ b/kernels/sgemm/sgemm.cu
@@ -15,8 +15,8 @@
 #define FLOAT4(value) (reinterpret_cast<float4 *>(&(value))[0])
 
 // modified from: https://zhuanlan.zhihu.com/p/657632577
-//  FP32
-//  SGEMM naive: compute one c[i,j]
+// FP32
+// SGEMM naive: compute one c[i,j]
 // element per threads, all row major
 __global__ void sgemm_naive_f32_kernel(float *a, float *b, float *c, int M,
                                        int N, int K) {
@@ -496,8 +496,6 @@ __global__ void sgemm_t_8x8_sliced_k_f32x4_bcf_dbuf_kernel(
   }
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/sgemm/sgemm_async.cu b/kernels/sgemm/sgemm_async.cu
index 2e562adb..d19beb83 100644
--- a/kernels/sgemm/sgemm_async.cu
+++ b/kernels/sgemm/sgemm_async.cu
@@ -812,8 +812,6 @@ __global__ void sgemm_t_8x16_sliced_k16_f32x4_bcf_dbuf_async_kernel(
   }
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/sgemm/sgemm_wmma_tf32_stage.cu b/kernels/sgemm/sgemm_wmma_tf32_stage.cu
index 09c10bba..503c3ec1 100644
--- a/kernels/sgemm/sgemm_wmma_tf32_stage.cu
+++ b/kernels/sgemm/sgemm_wmma_tf32_stage.cu
@@ -484,8 +484,6 @@ __global__ void sgemm_wmma_m16n16k8_mma4x2_warp2x4_stages_dsmem_kernel(
   }
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/sigmoid/sigmoid.cu b/kernels/sigmoid/sigmoid.cu
index 8bab212d..d5c8bf35 100644
--- a/kernels/sigmoid/sigmoid.cu
+++ b/kernels/sigmoid/sigmoid.cu
@@ -21,8 +21,8 @@
 #define MAX_EXP_F16 __float2half(11.089866488461016f)
 #define MIN_EXP_F16 __float2half(-9.704060527839234f)
 
-//  FP32
-//  Sigmoid x: N, y: N y=1/(1+exp(-x))
+// FP32
+// Sigmoid x: N, y: N y=1/(1+exp(-x))
 // grid(N/256), block(K=256)
 __global__ void sigmoid_f32_kernel(float *x, float *y, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
@@ -146,8 +146,6 @@ __global__ void sigmoid_f16x8_pack_kernel(half *x, half *y, int N) {
   }
 }
 
-// --------------------- PyTorch bindings for custom kernel
-// -----------------------
 #define STRINGFY(str) #str
 #define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
   m.def(STRINGFY(func), &func, STRINGFY(func));
diff --git a/kernels/swish/swish.cu b/kernels/swish/swish.cu
index d27a9afb..fb42efd1 100644
--- a/kernels/swish/swish.cu
+++ b/kernels/swish/swish.cu
@@ -15,8 +15,8 @@
 #define BFLOAT2(value) (reinterpret_cast<__nv_bfloat162 *>(&(value))[0])
 #define LDST128BITS(value) (reinterpret_cast<float4 *>(&(value))[0])
 
-//  FP32
-//  Swish x: N, y: N y=x*sigmoid(x)
+// FP32
+// Swish x: N, y: N y=x*sigmoid(x)
 __device__ __forceinline__ float swish(float x) {
   return x / (1.0f + expf(-x));
 }

From 76d28671c0b8f2614f99b8ae27ca78e9cf9acafc Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Sat, 28 Jun 2025 20:34:06 +0800
Subject: [PATCH 37/57] feat: triton layer-norm (#351)

* fix comments

* fix comments

* fix comments

* fix comments

* add torch.compile blogs

* add torch.compile blogs

* add torch.compile blogs

* feat: triton layer-norm

* feat: triton layer-norm

* feat: triton layer-norm

* feat: triton layer-norm

* feat: triton layer-norm

* feat: triton layer-norm

* feat: triton layer-norm

* Update README.md

* Update README.md
---
 README.md                                     |   8 +-
 .../__grp__add_kernel.json                    |   1 -
 .../add_kernel.cubin                          | Bin 6888 -> 0 bytes
 .../add_kernel.json                           |   1 -
 .../add_kernel.llir                           | 115 -----
 .../add_kernel.ptx                            | 272 ------------
 .../add_kernel.ttgir                          |  39 --
 .../add_kernel.ttir                           |  38 --
 .../{ffpa-attn => fused-attention}/.gitignore |   0
 kernels/openai-triton/hgemm/.gitignore        |  20 -
 .../{flash-attn => layer-norm}/.gitignore     |   0
 kernels/openai-triton/layer-norm/README.md    | 269 ++++++++++++
 kernels/openai-triton/layer-norm/bwd-math.png | Bin 0 -> 14307 bytes
 kernels/openai-triton/layer-norm/bwd.png      | Bin 0 -> 36018 bytes
 .../layer-norm/layer-norm-backward.csv        |  31 ++
 .../layer-norm/layer-norm-backward.png        | Bin 0 -> 32560 bytes
 .../layer-norm/layer-norm-forward.csv         |  31 ++
 .../layer-norm/layer-norm-forward.png         | Bin 0 -> 50694 bytes
 kernels/openai-triton/layer-norm/results.html |   3 +
 .../layer-norm/triton_layer_norm.py           | 401 ++++++++++++++++++
 .../.gitignore                                |   0
 kernels/openai-triton/reduce/.gitignore       |  20 -
 kernels/openai-triton/sgemm/.gitignore        |  20 -
 .../{elementwise => vector-add}/.gitignore    |   0
 .../{elementwise => vector-add}/README.md     |   0
 .../{elementwise => vector-add}/results.html  |   0
 .../triton_vector_add.py                      |   0
 .../vector-add-performance.png                | Bin
 28 files changed, 740 insertions(+), 529 deletions(-)
 delete mode 100644 kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/__grp__add_kernel.json
 delete mode 100644 kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.cubin
 delete mode 100644 kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.json
 delete mode 100644 kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.llir
 delete mode 100644 kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.ptx
 delete mode 100644 kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.ttgir
 delete mode 100644 kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.ttir
 rename kernels/openai-triton/{ffpa-attn => fused-attention}/.gitignore (100%)
 delete mode 100644 kernels/openai-triton/hgemm/.gitignore
 rename kernels/openai-triton/{flash-attn => layer-norm}/.gitignore (100%)
 create mode 100644 kernels/openai-triton/layer-norm/README.md
 create mode 100644 kernels/openai-triton/layer-norm/bwd-math.png
 create mode 100644 kernels/openai-triton/layer-norm/bwd.png
 create mode 100644 kernels/openai-triton/layer-norm/layer-norm-backward.csv
 create mode 100644 kernels/openai-triton/layer-norm/layer-norm-backward.png
 create mode 100644 kernels/openai-triton/layer-norm/layer-norm-forward.csv
 create mode 100644 kernels/openai-triton/layer-norm/layer-norm-forward.png
 create mode 100644 kernels/openai-triton/layer-norm/results.html
 create mode 100644 kernels/openai-triton/layer-norm/triton_layer_norm.py
 rename kernels/openai-triton/{flash-mla => matrix-multiplication}/.gitignore (100%)
 delete mode 100644 kernels/openai-triton/reduce/.gitignore
 delete mode 100644 kernels/openai-triton/sgemm/.gitignore
 rename kernels/openai-triton/{elementwise => vector-add}/.gitignore (100%)
 rename kernels/openai-triton/{elementwise => vector-add}/README.md (100%)
 rename kernels/openai-triton/{elementwise => vector-add}/results.html (100%)
 rename kernels/openai-triton/{elementwise => vector-add}/triton_vector_add.py (100%)
 rename kernels/openai-triton/{elementwise => vector-add}/vector-add-performance.png (100%)

diff --git a/README.md b/README.md
index d97af656..c7451416 100644
--- a/README.md
+++ b/README.md
@@ -10,7 +10,7 @@
       <img src=https://img.shields.io/github/watchers/xlite-dev/LeetCUDA?color=9cc >
       <img src=https://img.shields.io/github/forks/xlite-dev/LeetCUDA.svg?style=social >
       <img src=https://img.shields.io/github/stars/xlite-dev/LeetCUDA.svg?style=social >
-      <img src=https://img.shields.io/badge/Release-v3.0.6-brightgreen.svg >
+      <img src=https://img.shields.io/badge/Release-v3.0.12-brightgreen.svg >
       <img src=https://img.shields.io/badge/License-GPLv3.0-turquoise.svg >
   </div>
 </div>
@@ -459,8 +459,10 @@ The kernels listed here will guide you through a step-by-step progression, rangi
 
 |📖 Triton Kernel| 📖 Elem DType| 📖 Acc DType| 📖 Docs | 📖 Level |
 |:---|:---|:---|:---|:---|
-| ✔️ [triton_vector_add_kernel](./kernels/openai-triton/elementwise/)|all|all|[link](./kernels/openai-triton/elementwise/)|⭐️⭐️|
-| ✔️ [triton_fused_softmax(multi-stages)](./kernels/openai-triton/fused-softmax/)|f16/bf16/f32|f32|[link](./kernels/openai-triton/fused-softmax//)|⭐️⭐️⭐️|
+| ✔️ [triton_vector_add_kernel](./kernels/openai-triton/vector-add/)|all|all|[link](./kernels/openai-triton/vector-add/)|⭐️⭐️|
+| ✔️ [triton_fused_softmax(multi-stages)](./kernels/openai-triton/fused-softmax/)|f16/bf16/f32|f32|[link](./kernels/openai-triton/fused-softmax/)|⭐️⭐️⭐️|
+| ✔️ [triton_fused_layer_norm(forward-pass)](./kernels/openai-triton/layer-norm/)|f16/bf16/f32|f32|[link](./kernels/openai-triton/layer-norm/)|⭐️⭐️⭐️|
+| ✔️ [triton_fused_layer_norm(backward-pass)](./kernels/openai-triton/layer-norm/)|f16/bf16/f32|f32|[link](./kernels/openai-triton/layer-norm/)|⭐️⭐️⭐️|
 | ✔️ [triton_merge_attn_states_kernel(w/ CUDA)](./kernels/openai-triton/merge-attn-states/)|f16/bf16/f32|f32|[link](./kernels/openai-triton/merge-attn-states/)|⭐️⭐️⭐️|
 
 ### 📚 CUTLASS/CuTe Kernel ⭐️⭐️⭐️ ([©️back👆🏻](#cuda-kernel))
diff --git a/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/__grp__add_kernel.json b/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/__grp__add_kernel.json
deleted file mode 100644
index fd1fc70b..00000000
--- a/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/__grp__add_kernel.json
+++ /dev/null
@@ -1 +0,0 @@
-{"child_paths": {"add_kernel.ttir": "/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.ttir", "add_kernel.ttgir": "/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.ttgir", "add_kernel.llir": "/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.llir", "add_kernel.ptx": "/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.ptx", "add_kernel.cubin": "/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.cubin", "add_kernel.json": "/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.json"}}
diff --git a/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.cubin b/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.cubin
deleted file mode 100644
index 3ed48bb4459a5abd6ecbe9493860979c728883a0..0000000000000000000000000000000000000000
GIT binary patch
diff --git a/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.json b/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.json
deleted file mode 100644
index 85013cd9..00000000
--- a/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.json
+++ /dev/null
@@ -1 +0,0 @@
-{"hash": "c8228ac8c26b356b0eaa146a6aa48d1cfa27192224300c73e65ad0119f459db4", "target": {"backend": "cuda", "arch": 89, "warp_size": 32}, "num_warps": 4, "num_ctas": 1, "num_stages": 3, "maxnreg": null, "cluster_dims": [1, 1, 1], "ptx_version": null, "enable_fp_fusion": true, "launch_cooperative_grid": false, "supported_fp8_dtypes": ["fp8e4b15", "fp8e4nv", "fp8e5"], "deprecated_fp8_dtypes": [], "default_dot_input_precision": "tf32", "allowed_dot_input_precisions": ["tf32", "tf32x3", "ieee"], "max_num_imprecise_acc_default": 0, "extern_libs": [["libdevice", "/root/anaconda3/envs/triton-nightly/lib/python3.10/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]], "debug": false, "backend_name": "cuda", "sanitize_overflow": true, "arch": "sm89", "triton_version": "3.3.0", "shared": 0, "tmem_size": 0, "global_scratch_size": 0, "global_scratch_align": 1, "name": "add_kernel"}
diff --git a/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.llir b/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.llir
deleted file mode 100644
index 7ec7b9d6..00000000
--- a/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.llir
+++ /dev/null
@@ -1,115 +0,0 @@
-; ModuleID = 'LLVMDialectModule'
-source_filename = "LLVMDialectModule"
-target datalayout = "e-p6:32:32-i64:64-i128:128-v16:16-v32:32-n16:32:64"
-
-define ptx_kernel void @add_kernel(ptr addrspace(1) %0, ptr addrspace(1) %1, ptr addrspace(1) %2, i32 %3, ptr addrspace(1) readnone captures(none) %4) local_unnamed_addr !dbg !6 {
-  %6 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !9
-  %7 = shl i32 %6, 10, !dbg !10
-  %8 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !11
-  %9 = shl i32 %8, 2, !dbg !11
-  %10 = and i32 %9, 508, !dbg !11
-  %11 = or disjoint i32 %10, %7, !dbg !12
-  %12 = or disjoint i32 %11, 512, !dbg !12
-  %13 = icmp slt i32 %11, %3, !dbg !13
-  %14 = icmp slt i32 %12, %3, !dbg !13
-  %15 = sext i32 %11 to i64, !dbg !14
-  %16 = getelementptr float, ptr addrspace(1) %0, i64 %15, !dbg !14
-  %17 = sext i32 %12 to i64, !dbg !14
-  %18 = getelementptr float, ptr addrspace(1) %0, i64 %17, !dbg !14
-  %19 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];", "=r,=r,=r,=r,l,b"(ptr addrspace(1) %16, i1 %13) #1, !dbg !15
-  %20 = extractvalue { i32, i32, i32, i32 } %19, 0, !dbg !15
-  %21 = extractvalue { i32, i32, i32, i32 } %19, 1, !dbg !15
-  %22 = extractvalue { i32, i32, i32, i32 } %19, 2, !dbg !15
-  %23 = extractvalue { i32, i32, i32, i32 } %19, 3, !dbg !15
-  %24 = bitcast i32 %20 to float, !dbg !15
-  %25 = bitcast i32 %21 to float, !dbg !15
-  %26 = bitcast i32 %22 to float, !dbg !15
-  %27 = bitcast i32 %23 to float, !dbg !15
-  %28 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];", "=r,=r,=r,=r,l,b"(ptr addrspace(1) %18, i1 %14) #1, !dbg !15
-  %29 = extractvalue { i32, i32, i32, i32 } %28, 0, !dbg !15
-  %30 = extractvalue { i32, i32, i32, i32 } %28, 1, !dbg !15
-  %31 = extractvalue { i32, i32, i32, i32 } %28, 2, !dbg !15
-  %32 = extractvalue { i32, i32, i32, i32 } %28, 3, !dbg !15
-  %33 = bitcast i32 %29 to float, !dbg !15
-  %34 = bitcast i32 %30 to float, !dbg !15
-  %35 = bitcast i32 %31 to float, !dbg !15
-  %36 = bitcast i32 %32 to float, !dbg !15
-  %37 = getelementptr float, ptr addrspace(1) %1, i64 %15, !dbg !16
-  %38 = getelementptr float, ptr addrspace(1) %1, i64 %17, !dbg !16
-  %39 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];", "=r,=r,=r,=r,l,b"(ptr addrspace(1) %37, i1 %13) #1, !dbg !17
-  %40 = extractvalue { i32, i32, i32, i32 } %39, 0, !dbg !17
-  %41 = extractvalue { i32, i32, i32, i32 } %39, 1, !dbg !17
-  %42 = extractvalue { i32, i32, i32, i32 } %39, 2, !dbg !17
-  %43 = extractvalue { i32, i32, i32, i32 } %39, 3, !dbg !17
-  %44 = bitcast i32 %40 to float, !dbg !17
-  %45 = bitcast i32 %41 to float, !dbg !17
-  %46 = bitcast i32 %42 to float, !dbg !17
-  %47 = bitcast i32 %43 to float, !dbg !17
-  %48 = tail call { i32, i32, i32, i32 } asm sideeffect "mov.u32 $0, 0x0;\0A\09mov.u32 $1, 0x0;\0A\09mov.u32 $2, 0x0;\0A\09mov.u32 $3, 0x0;\0A\09@$5 ld.global.v4.b32 { $0, $1, $2, $3 }, [ $4 + 0 ];", "=r,=r,=r,=r,l,b"(ptr addrspace(1) %38, i1 %14) #1, !dbg !17
-  %49 = extractvalue { i32, i32, i32, i32 } %48, 0, !dbg !17
-  %50 = extractvalue { i32, i32, i32, i32 } %48, 1, !dbg !17
-  %51 = extractvalue { i32, i32, i32, i32 } %48, 2, !dbg !17
-  %52 = extractvalue { i32, i32, i32, i32 } %48, 3, !dbg !17
-  %53 = bitcast i32 %49 to float, !dbg !17
-  %54 = bitcast i32 %50 to float, !dbg !17
-  %55 = bitcast i32 %51 to float, !dbg !17
-  %56 = bitcast i32 %52 to float, !dbg !17
-  %57 = fadd float %24, %44, !dbg !18
-  %58 = fadd float %25, %45, !dbg !18
-  %59 = fadd float %26, %46, !dbg !18
-  %60 = fadd float %27, %47, !dbg !18
-  %61 = fadd float %33, %53, !dbg !18
-  %62 = fadd float %34, %54, !dbg !18
-  %63 = fadd float %35, %55, !dbg !18
-  %64 = fadd float %36, %56, !dbg !18
-  %65 = getelementptr float, ptr addrspace(1) %2, i64 %15, !dbg !19
-  %66 = getelementptr float, ptr addrspace(1) %2, i64 %17, !dbg !19
-  %67 = bitcast float %57 to i32, !dbg !20
-  %68 = bitcast float %58 to i32, !dbg !20
-  %69 = bitcast float %59 to i32, !dbg !20
-  %70 = bitcast float %60 to i32, !dbg !20
-  tail call void asm sideeffect "@$5 st.global.v4.b32 [ $4 + 0 ], { $0, $1, $2, $3 };", "r,r,r,r,l,b"(i32 %67, i32 %68, i32 %69, i32 %70, ptr addrspace(1) %65, i1 %13) #1, !dbg !20
-  %71 = bitcast float %61 to i32, !dbg !20
-  %72 = bitcast float %62 to i32, !dbg !20
-  %73 = bitcast float %63 to i32, !dbg !20
-  %74 = bitcast float %64 to i32, !dbg !20
-  tail call void asm sideeffect "@$5 st.global.v4.b32 [ $4 + 0 ], { $0, $1, $2, $3 };", "r,r,r,r,l,b"(i32 %71, i32 %72, i32 %73, i32 %74, ptr addrspace(1) %66, i1 %14) #1, !dbg !20
-  ret void, !dbg !21
-}
-
-; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
-declare noundef i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #0
-
-; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
-declare noundef i32 @llvm.nvvm.read.ptx.sreg.tid.x() #0
-
-attributes #0 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }
-attributes #1 = { nounwind }
-
-!llvm.module.flags = !{!0, !1}
-!llvm.dbg.cu = !{!2}
-!nvvm.annotations = !{!4}
-!llvm.ident = !{!5}
-
-!0 = !{i32 2, !"Debug Info Version", i32 3}
-!1 = !{i32 4, !"nvvm-reflect-ftz", i32 1}
-!2 = distinct !DICompileUnit(language: DW_LANG_C, file: !3, producer: "triton", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)
-!3 = !DIFile(filename: "triton_vector_add.py", directory: "/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise")
-!4 = !{ptr @add_kernel, !"reqntidx", i32 128}
-!5 = !{!"clang version 3.8.0 (tags/RELEASE_380/final)"}
-!6 = distinct !DISubprogram(name: "add_kernel", linkageName: "add_kernel", scope: !3, file: !3, line: 32, type: !7, scopeLine: 32, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !2)
-!7 = !DISubroutineType(cc: DW_CC_normal, types: !8)
-!8 = !{}
-!9 = !DILocation(line: 41, column: 24, scope: !6)
-!10 = !DILocation(line: 46, column: 24, scope: !6)
-!11 = !DILocation(line: 47, column: 41, scope: !6)
-!12 = !DILocation(line: 47, column: 28, scope: !6)
-!13 = !DILocation(line: 49, column: 21, scope: !6)
-!14 = !DILocation(line: 52, column: 24, scope: !6)
-!15 = !DILocation(line: 52, column: 16, scope: !6)
-!16 = !DILocation(line: 53, column: 24, scope: !6)
-!17 = !DILocation(line: 53, column: 16, scope: !6)
-!18 = !DILocation(line: 54, column: 17, scope: !6)
-!19 = !DILocation(line: 56, column: 26, scope: !6)
-!20 = !DILocation(line: 56, column: 35, scope: !6)
-!21 = !DILocation(line: 56, column: 4, scope: !6)
diff --git a/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.ptx b/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.ptx
deleted file mode 100644
index 53545e03..00000000
--- a/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.ptx
+++ /dev/null
@@ -1,272 +0,0 @@
-//
-// Generated by LLVM NVPTX Back-End
-//
-
-.version 8.4
-.target sm_89
-.address_size 64
-
-	// .globl	add_kernel              // -- Begin function add_kernel
-                                        // @add_kernel
-.visible .entry add_kernel(
-	.param .u64 .ptr .global .align 1 add_kernel_param_0,
-	.param .u64 .ptr .global .align 1 add_kernel_param_1,
-	.param .u64 .ptr .global .align 1 add_kernel_param_2,
-	.param .u32 add_kernel_param_3,
-	.param .u64 .ptr .global .align 1 add_kernel_param_4
-)
-.reqntid 128, 1, 1
-{
-	.reg .pred 	%p<7>;
-	.reg .b32 	%r<33>;
-	.reg .f32 	%f<25>;
-	.reg .b64 	%rd<11>;
-	.loc	1 32 0                          // triton_vector_add.py:32:0
-$L__func_begin0:
-	.loc	1 32 0                          // triton_vector_add.py:32:0
-
-// %bb.0:
-	ld.param.u64 	%rd7, [add_kernel_param_0];
-	ld.param.u64 	%rd8, [add_kernel_param_1];
-$L__tmp0:
-	.loc	1 41 24                         // triton_vector_add.py:41:24
-	mov.u32 	%r25, %ctaid.x;
-	.loc	1 46 24                         // triton_vector_add.py:46:24
-	shl.b32 	%r26, %r25, 10;
-	ld.param.u64 	%rd9, [add_kernel_param_2];
-	ld.param.u32 	%r27, [add_kernel_param_3];
-	.loc	1 47 41                         // triton_vector_add.py:47:41
-	mov.u32 	%r28, %tid.x;
-	shl.b32 	%r29, %r28, 2;
-	and.b32  	%r30, %r29, 508;
-	.loc	1 47 28                         // triton_vector_add.py:47:28
-	or.b32  	%r31, %r30, %r26;
-	or.b32  	%r32, %r31, 512;
-	.loc	1 49 21                         // triton_vector_add.py:49:21
-	setp.lt.s32 	%p1, %r31, %r27;
-	setp.lt.s32 	%p2, %r32, %r27;
-	.loc	1 52 24                         // triton_vector_add.py:52:24
-	mul.wide.s32 	%rd10, %r31, 4;
-	add.s64 	%rd1, %rd7, %rd10;
-	add.s64 	%rd2, %rd1, 2048;
-	.loc	1 52 16                         // triton_vector_add.py:52:16
-	// begin inline asm
-	mov.u32 %r1, 0x0;
-	mov.u32 %r2, 0x0;
-	mov.u32 %r3, 0x0;
-	mov.u32 %r4, 0x0;
-	@%p1 ld.global.v4.b32 { %r1, %r2, %r3, %r4 }, [ %rd1 + 0 ];
-	// end inline asm
-	mov.b32 	%f1, %r1;
-	mov.b32 	%f2, %r2;
-	mov.b32 	%f3, %r3;
-	mov.b32 	%f4, %r4;
-	// begin inline asm
-	mov.u32 %r5, 0x0;
-	mov.u32 %r6, 0x0;
-	mov.u32 %r7, 0x0;
-	mov.u32 %r8, 0x0;
-	@%p2 ld.global.v4.b32 { %r5, %r6, %r7, %r8 }, [ %rd2 + 0 ];
-	// end inline asm
-	mov.b32 	%f5, %r5;
-	mov.b32 	%f6, %r6;
-	mov.b32 	%f7, %r7;
-	mov.b32 	%f8, %r8;
-	.loc	1 53 24                         // triton_vector_add.py:53:24
-	add.s64 	%rd3, %rd8, %rd10;
-	add.s64 	%rd4, %rd3, 2048;
-	.loc	1 53 16                         // triton_vector_add.py:53:16
-	// begin inline asm
-	mov.u32 %r9, 0x0;
-	mov.u32 %r10, 0x0;
-	mov.u32 %r11, 0x0;
-	mov.u32 %r12, 0x0;
-	@%p1 ld.global.v4.b32 { %r9, %r10, %r11, %r12 }, [ %rd3 + 0 ];
-	// end inline asm
-	mov.b32 	%f9, %r9;
-	mov.b32 	%f10, %r10;
-	mov.b32 	%f11, %r11;
-	mov.b32 	%f12, %r12;
-	// begin inline asm
-	mov.u32 %r13, 0x0;
-	mov.u32 %r14, 0x0;
-	mov.u32 %r15, 0x0;
-	mov.u32 %r16, 0x0;
-	@%p2 ld.global.v4.b32 { %r13, %r14, %r15, %r16 }, [ %rd4 + 0 ];
-	// end inline asm
-	mov.b32 	%f13, %r13;
-	mov.b32 	%f14, %r14;
-	mov.b32 	%f15, %r15;
-	mov.b32 	%f16, %r16;
-	.loc	1 54 17                         // triton_vector_add.py:54:17
-	add.f32 	%f17, %f1, %f9;
-	add.f32 	%f18, %f2, %f10;
-	add.f32 	%f19, %f3, %f11;
-	add.f32 	%f20, %f4, %f12;
-	add.f32 	%f21, %f5, %f13;
-	add.f32 	%f22, %f6, %f14;
-	add.f32 	%f23, %f7, %f15;
-	add.f32 	%f24, %f8, %f16;
-	.loc	1 56 26                         // triton_vector_add.py:56:26
-	add.s64 	%rd5, %rd9, %rd10;
-	add.s64 	%rd6, %rd5, 2048;
-	.loc	1 56 35                         // triton_vector_add.py:56:35
-	mov.b32 	%r17, %f17;
-	mov.b32 	%r18, %f18;
-	mov.b32 	%r19, %f19;
-	mov.b32 	%r20, %f20;
-	// begin inline asm
-	@%p1 st.global.v4.b32 [ %rd5 + 0 ], { %r17, %r18, %r19, %r20 };
-	// end inline asm
-	mov.b32 	%r21, %f21;
-	mov.b32 	%r22, %f22;
-	mov.b32 	%r23, %f23;
-	mov.b32 	%r24, %f24;
-	// begin inline asm
-	@%p2 st.global.v4.b32 [ %rd6 + 0 ], { %r21, %r22, %r23, %r24 };
-	// end inline asm
-	.loc	1 56 4                          // triton_vector_add.py:56:4
-	ret;
-$L__tmp1:
-$L__func_end0:
-                                        // -- End function
-}
-	.file	1 "/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py"
-	.section	.debug_abbrev
-	{
-.b8 1                                   // Abbreviation Code
-.b8 17                                  // DW_TAG_compile_unit
-.b8 0                                   // DW_CHILDREN_no
-.b8 37                                  // DW_AT_producer
-.b8 8                                   // DW_FORM_string
-.b8 19                                  // DW_AT_language
-.b8 5                                   // DW_FORM_data2
-.b8 3                                   // DW_AT_name
-.b8 8                                   // DW_FORM_string
-.b8 16                                  // DW_AT_stmt_list
-.b8 6                                   // DW_FORM_data4
-.b8 27                                  // DW_AT_comp_dir
-.b8 8                                   // DW_FORM_string
-.b8 0                                   // EOM(1)
-.b8 0                                   // EOM(2)
-.b8 0                                   // EOM(3)
-	}
-	.section	.debug_info
-	{
-.b32 120                                // Length of Unit
-.b8 2                                   // DWARF version number
-.b8 0
-.b32 .debug_abbrev                      // Offset Into Abbrev. Section
-.b8 8                                   // Address Size (in bytes)
-.b8 1                                   // Abbrev [1] 0xb:0x71 DW_TAG_compile_unit
-.b8 116                                 // DW_AT_producer
-.b8 114
-.b8 105
-.b8 116
-.b8 111
-.b8 110
-.b8 0
-.b8 2                                   // DW_AT_language
-.b8 0
-.b8 116                                 // DW_AT_name
-.b8 114
-.b8 105
-.b8 116
-.b8 111
-.b8 110
-.b8 95
-.b8 118
-.b8 101
-.b8 99
-.b8 116
-.b8 111
-.b8 114
-.b8 95
-.b8 97
-.b8 100
-.b8 100
-.b8 46
-.b8 112
-.b8 121
-.b8 0
-.b32 .debug_line                        // DW_AT_stmt_list
-.b8 47                                  // DW_AT_comp_dir
-.b8 119
-.b8 111
-.b8 114
-.b8 107
-.b8 115
-.b8 112
-.b8 97
-.b8 99
-.b8 101
-.b8 47
-.b8 100
-.b8 101
-.b8 118
-.b8 47
-.b8 111
-.b8 112
-.b8 101
-.b8 110
-.b8 108
-.b8 108
-.b8 109
-.b8 47
-.b8 100
-.b8 101
-.b8 118
-.b8 47
-.b8 67
-.b8 85
-.b8 68
-.b8 65
-.b8 45
-.b8 76
-.b8 101
-.b8 97
-.b8 114
-.b8 110
-.b8 45
-.b8 78
-.b8 111
-.b8 116
-.b8 101
-.b8 115
-.b8 47
-.b8 107
-.b8 101
-.b8 114
-.b8 110
-.b8 101
-.b8 108
-.b8 115
-.b8 47
-.b8 111
-.b8 112
-.b8 101
-.b8 110
-.b8 97
-.b8 105
-.b8 45
-.b8 116
-.b8 114
-.b8 105
-.b8 116
-.b8 111
-.b8 110
-.b8 47
-.b8 101
-.b8 108
-.b8 101
-.b8 109
-.b8 101
-.b8 110
-.b8 116
-.b8 119
-.b8 105
-.b8 115
-.b8 101
-.b8 0
-	}
-	.section	.debug_macinfo	{	}
diff --git a/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.ttgir b/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.ttgir
deleted file mode 100644
index 7704d23c..00000000
--- a/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.ttgir
+++ /dev/null
@@ -1,39 +0,0 @@
-#blocked = #ttg.blocked<{sizePerThread = [4], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>
-#loc = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":32:0)
-module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.target = "cuda:89", "ttg.threads-per-warp" = 32 : i32} {
-  tt.func public @add_kernel(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":32:0), %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":32:0), %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":32:0), %arg3: i32 {tt.divisibility = 16 : i32} loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":32:0)) attributes {noinline = false} {
-    %c1024_i32 = arith.constant 1024 : i32 loc(#loc1)
-    %0 = tt.get_program_id x : i32 loc(#loc2)
-    %1 = arith.muli %0, %c1024_i32 : i32 loc(#loc3)
-    %2 = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32, #blocked> loc(#loc4)
-    %3 = tt.splat %1 : i32 -> tensor<1024xi32, #blocked> loc(#loc5)
-    %4 = arith.addi %3, %2 : tensor<1024xi32, #blocked> loc(#loc5)
-    %5 = tt.splat %arg3 : i32 -> tensor<1024xi32, #blocked> loc(#loc6)
-    %6 = arith.cmpi slt, %4, %5 : tensor<1024xi32, #blocked> loc(#loc6)
-    %7 = tt.splat %arg0 : !tt.ptr<f32> -> tensor<1024x!tt.ptr<f32>, #blocked> loc(#loc7)
-    %8 = tt.addptr %7, %4 : tensor<1024x!tt.ptr<f32>, #blocked>, tensor<1024xi32, #blocked> loc(#loc7)
-    %9 = tt.load %8, %6 : tensor<1024x!tt.ptr<f32>, #blocked> loc(#loc8)
-    %10 = tt.splat %arg1 : !tt.ptr<f32> -> tensor<1024x!tt.ptr<f32>, #blocked> loc(#loc9)
-    %11 = tt.addptr %10, %4 : tensor<1024x!tt.ptr<f32>, #blocked>, tensor<1024xi32, #blocked> loc(#loc9)
-    %12 = tt.load %11, %6 : tensor<1024x!tt.ptr<f32>, #blocked> loc(#loc10)
-    %13 = arith.addf %9, %12 : tensor<1024xf32, #blocked> loc(#loc11)
-    %14 = tt.splat %arg2 : !tt.ptr<f32> -> tensor<1024x!tt.ptr<f32>, #blocked> loc(#loc12)
-    %15 = tt.addptr %14, %4 : tensor<1024x!tt.ptr<f32>, #blocked>, tensor<1024xi32, #blocked> loc(#loc12)
-    tt.store %15, %13, %6 : tensor<1024x!tt.ptr<f32>, #blocked> loc(#loc13)
-    tt.return loc(#loc14)
-  } loc(#loc)
-} loc(#loc)
-#loc1 = loc(unknown)
-#loc2 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":41:24)
-#loc3 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":46:24)
-#loc4 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":47:41)
-#loc5 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":47:28)
-#loc6 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":49:21)
-#loc7 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":52:24)
-#loc8 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":52:16)
-#loc9 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":53:24)
-#loc10 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":53:16)
-#loc11 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":54:17)
-#loc12 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":56:26)
-#loc13 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":56:35)
-#loc14 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":56:4)
diff --git a/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.ttir b/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.ttir
deleted file mode 100644
index e45561f5..00000000
--- a/kernels/openai-triton/elementwise/cache/ZARIVSGCNM2WWDVKCRVGVJENDT5COGJCEQYAY47GLLIBDH2FTW2A/add_kernel.ttir
+++ /dev/null
@@ -1,38 +0,0 @@
-#loc = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":32:0)
-module {
-  tt.func public @add_kernel(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":32:0), %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":32:0), %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":32:0), %arg3: i32 {tt.divisibility = 16 : i32} loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":32:0)) attributes {noinline = false} {
-    %c1024_i32 = arith.constant 1024 : i32 loc(#loc1)
-    %0 = tt.get_program_id x : i32 loc(#loc2)
-    %1 = arith.muli %0, %c1024_i32 : i32 loc(#loc3)
-    %2 = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32> loc(#loc4)
-    %3 = tt.splat %1 : i32 -> tensor<1024xi32> loc(#loc5)
-    %4 = arith.addi %3, %2 : tensor<1024xi32> loc(#loc5)
-    %5 = tt.splat %arg3 : i32 -> tensor<1024xi32> loc(#loc6)
-    %6 = arith.cmpi slt, %4, %5 : tensor<1024xi32> loc(#loc6)
-    %7 = tt.splat %arg0 : !tt.ptr<f32> -> tensor<1024x!tt.ptr<f32>> loc(#loc7)
-    %8 = tt.addptr %7, %4 : tensor<1024x!tt.ptr<f32>>, tensor<1024xi32> loc(#loc7)
-    %9 = tt.load %8, %6 : tensor<1024x!tt.ptr<f32>> loc(#loc8)
-    %10 = tt.splat %arg1 : !tt.ptr<f32> -> tensor<1024x!tt.ptr<f32>> loc(#loc9)
-    %11 = tt.addptr %10, %4 : tensor<1024x!tt.ptr<f32>>, tensor<1024xi32> loc(#loc9)
-    %12 = tt.load %11, %6 : tensor<1024x!tt.ptr<f32>> loc(#loc10)
-    %13 = arith.addf %9, %12 : tensor<1024xf32> loc(#loc11)
-    %14 = tt.splat %arg2 : !tt.ptr<f32> -> tensor<1024x!tt.ptr<f32>> loc(#loc12)
-    %15 = tt.addptr %14, %4 : tensor<1024x!tt.ptr<f32>>, tensor<1024xi32> loc(#loc12)
-    tt.store %15, %13, %6 : tensor<1024x!tt.ptr<f32>> loc(#loc13)
-    tt.return loc(#loc14)
-  } loc(#loc)
-} loc(#loc)
-#loc1 = loc(unknown)
-#loc2 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":41:24)
-#loc3 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":46:24)
-#loc4 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":47:41)
-#loc5 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":47:28)
-#loc6 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":49:21)
-#loc7 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":52:24)
-#loc8 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":52:16)
-#loc9 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":53:24)
-#loc10 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":53:16)
-#loc11 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":54:17)
-#loc12 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":56:26)
-#loc13 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":56:35)
-#loc14 = loc("/workspace/dev/openllm/dev/CUDA-Learn-Notes/kernels/openai-triton/elementwise/triton_vector_add.py":56:4)
diff --git a/kernels/openai-triton/ffpa-attn/.gitignore b/kernels/openai-triton/fused-attention/.gitignore
similarity index 100%
rename from kernels/openai-triton/ffpa-attn/.gitignore
rename to kernels/openai-triton/fused-attention/.gitignore
diff --git a/kernels/openai-triton/hgemm/.gitignore b/kernels/openai-triton/hgemm/.gitignore
deleted file mode 100644
index a9886a40..00000000
--- a/kernels/openai-triton/hgemm/.gitignore
+++ /dev/null
@@ -1,20 +0,0 @@
-*.so
-*.a
-*.dylib
-*.dll
-*.lib
-.DS_Store
-build
-*.whl
-tmp
-*.onnx
-*.pb
-*.pbtxt
-*.pt
-*.pth
-*.engine
-*.bin
-*.nsys
-*.nvvp
-*.nsys*
-*.sqlite
diff --git a/kernels/openai-triton/flash-attn/.gitignore b/kernels/openai-triton/layer-norm/.gitignore
similarity index 100%
rename from kernels/openai-triton/flash-attn/.gitignore
rename to kernels/openai-triton/layer-norm/.gitignore
diff --git a/kernels/openai-triton/layer-norm/README.md b/kernels/openai-triton/layer-norm/README.md
new file mode 100644
index 00000000..4977f619
--- /dev/null
+++ b/kernels/openai-triton/layer-norm/README.md
@@ -0,0 +1,269 @@
+# Fused Layer Norm
+
+## Forward pass
+
+The *LayerNorm* operator was first introduced is a way to improve the performance of sequential models (e.g., Transformers) or neural networks with small batch size. It takes a vector $x$ as input and produces a vector $y$ of the same shape as output. The normalization is performed by subtracting the mean and dividing by the standard deviation of $x$. After the normalization, a learnable linear transformation with weights $w$ and biases $b$ is applied. The forward pass can be expressed as follows:
+
+$$y = \frac{ x - \text{E}[x] }{ \sqrt{\text{Var}(x) + \epsilon} } * w + b$$
+
+where $\epsilon$ is a small constant added to the denominator for numerical stability.
+
+
+```python
+@triton.jit
+def _layer_norm_fwd_fused(
+    X,  # pointer to the input
+    Y,  # pointer to the output
+    W,  # pointer to the weights
+    B,  # pointer to the biases
+    Mean,  # pointer to the mean
+    Rstd,  # pointer to the 1/std
+    stride,  # how much to increase the pointer when moving by 1 row
+    N,  # number of columns in X
+    eps,  # epsilon to avoid division by zero
+    BLOCK_SIZE: tl.constexpr,
+):
+    # Map the program id to the row of X and Y it should compute.
+    row = tl.program_id(0)
+    Y += row * stride
+    X += row * stride
+    # Compute mean
+    mean = 0
+    _mean = tl.zeros([BLOCK_SIZE], dtype=tl.float32)
+    for off in range(0, N, BLOCK_SIZE):
+        cols = off + tl.arange(0, BLOCK_SIZE)
+        a = tl.load(X + cols, mask=cols < N, other=0.0).to(tl.float32)
+        _mean += a
+    mean = tl.sum(_mean, axis=0) / N
+    # Compute variance
+    _var = tl.zeros([BLOCK_SIZE], dtype=tl.float32)
+    for off in range(0, N, BLOCK_SIZE):
+        cols = off + tl.arange(0, BLOCK_SIZE)
+        x = tl.load(X + cols, mask=cols < N, other=0.0).to(tl.float32)
+        x = tl.where(cols < N, x - mean, 0.0)
+        _var += x * x
+    var = tl.sum(_var, axis=0) / N
+    rstd = 1 / tl.sqrt(var + eps)
+    # Write mean / rstd
+    tl.store(Mean + row, mean)
+    tl.store(Rstd + row, rstd)
+    # Normalize and apply linear transformation
+    for off in range(0, N, BLOCK_SIZE):
+        cols = off + tl.arange(0, BLOCK_SIZE)
+        mask = cols < N
+        w = tl.load(W + cols, mask=mask)
+        b = tl.load(B + cols, mask=mask)
+        x = tl.load(X + cols, mask=mask, other=0.0).to(tl.float32)
+        x_hat = (x - mean) * rstd
+        y = x_hat * w + b
+        # Write output
+        tl.store(Y + cols, y, mask=mask)
+```
+
+## Backward pass
+
+The backward pass for the layer normalization operator is a bit more involved than the forward pass. Let $\hat{x}$ be the normalized inputs $\frac{ x - \text{E}[x] }{ \sqrt{\text{Var}(x) + \epsilon} }$ before the linear transformation, the Vector-Jacobian Products (VJP) $\nabla_{x}$ of $x$ are given by:
+
+
+$$\nabla_{x}=\frac{1}{\sigma}\Big( \nabla_{y} \odot w - \underbrace{ \big( \frac{1}{N} \hat{x} \cdot (\nabla_{y} \odot w) \big) }\_{ c_1 } \odot \hat{x} - \underbrace{ \frac{1}{N} \nabla_{y} \cdot w }_{c_2} \Big)$$
+
+<!--
+![alt text](bwd-math.png)
+-->
+
+where $\odot$ denotes the element-wise multiplication, $\cdot$ denotes the dot product, and $\sigma$ is the standard deviation. $c_1$ and $c_2$ are intermediate constants that improve the readability of the following implementation. For the weights $w$ and biases $b$, the VJPs $\nabla_{w}$ and $\nabla_{b}$ are more straightforward:
+
+$$\nabla_{w} = \nabla_{y} \odot \hat{x} \quad \text{and} \quad \nabla_{b} = \nabla_{y}$$
+
+Since the same weights $w$ and biases $b$ are used for all rows in the same batch, their gradients need to sum up. To perform this step efficiently, we use a parallel reduction strategy: each kernel instance accumulates
+partial $\nabla_{w}$ and $\nabla_{b}$ across certain rows into one of `GROUP_SIZE_M` independent buffers. These buffers stay in the L2 cache and then are further reduced by another function to compute the actual $\nabla_{w}$ and $\nabla_{b}$.
+
+Let the number of input rows $M = 4$ and GROUP_SIZE_M = 2, here's a diagram of the parallel reduction strategy for $\nabla_{w}$ ($\nabla_{b}$ is omitted for brevity):
+
+![](./bwd.png)
+
+In Stage 1, the rows of X that have the same color share the same buffer and thus a lock is used to ensure that only one kernel instance writes to the buffer at a time. In Stage 2, the buffers are further reduced to compute the final $\nabla_{w}$ and $\nabla_{b}$. In the following implementation, Stage 1 is implemented by the function `_layer_norm_bwd_dx_fused` and Stage 2 is implemented by the function `_layer_norm_bwd_dwdb`
+
+```python
+@triton.jit
+def _layer_norm_bwd_dx_fused(
+    DX,  # pointer to the input gradient
+    DY,  # pointer to the output gradient
+    DW,  # pointer to the partial sum of weights gradient
+    DB,  # pointer to the partial sum of biases gradient
+    X,  # pointer to the input
+    W,  # pointer to the weights
+    Mean,  # pointer to the mean
+    Rstd,  # pointer to the 1/std
+    Lock,  # pointer to the lock
+    stride,  # how much to increase the pointer when moving by 1 row
+    N,  # number of columns in X
+    GROUP_SIZE_M: tl.constexpr,
+    BLOCK_SIZE_N: tl.constexpr,
+):
+    # Map the program id to the elements of X, DX, and DY it should compute.
+    row = tl.program_id(0)
+    cols = tl.arange(0, BLOCK_SIZE_N)
+    mask = cols < N
+    X += row * stride
+    DY += row * stride
+    DX += row * stride
+    # Offset locks and weights/biases gradient pointer for parallel reduction
+    lock_id = row % GROUP_SIZE_M
+    Lock += lock_id
+    Count = Lock + GROUP_SIZE_M
+    DW = DW + lock_id * N + cols
+    DB = DB + lock_id * N + cols
+    # Load data to SRAM
+    x = tl.load(X + cols, mask=mask, other=0).to(tl.float32)
+    dy = tl.load(DY + cols, mask=mask, other=0).to(tl.float32)
+    w = tl.load(W + cols, mask=mask).to(tl.float32)
+    mean = tl.load(Mean + row)
+    rstd = tl.load(Rstd + row)
+    # Compute dx
+    xhat = (x - mean) * rstd
+    wdy = w * dy
+    xhat = tl.where(mask, xhat, 0.0)
+    wdy = tl.where(mask, wdy, 0.0)
+    c1 = tl.sum(xhat * wdy, axis=0) / N
+    c2 = tl.sum(wdy, axis=0) / N
+    dx = (wdy - (xhat * c1 + c2)) * rstd
+    # Write dx
+    tl.store(DX + cols, dx, mask=mask)
+    # Accumulate partial sums for dw/db
+    partial_dw = (dy * xhat).to(w.dtype)
+    partial_db = (dy).to(w.dtype)
+    while tl.atomic_cas(Lock, 0, 1) == 1:
+        pass
+    count = tl.load(Count)
+    # First store doesn't accumulate
+    if count == 0:
+        tl.atomic_xchg(Count, 1)
+    else:
+        partial_dw += tl.load(DW, mask=mask)
+        partial_db += tl.load(DB, mask=mask)
+    tl.store(DW, partial_dw, mask=mask)
+    tl.store(DB, partial_db, mask=mask)
+
+    # need a barrier to ensure all threads finished before
+    # releasing the lock
+    tl.debug_barrier()
+
+    # Release the lock
+    tl.atomic_xchg(Lock, 0)
+
+
+@triton.jit
+def _layer_norm_bwd_dwdb(
+    DW,  # pointer to the partial sum of weights gradient
+    DB,  # pointer to the partial sum of biases gradient
+    FINAL_DW,  # pointer to the weights gradient
+    FINAL_DB,  # pointer to the biases gradient
+    M,  # GROUP_SIZE_M
+    N,  # number of columns
+    BLOCK_SIZE_M: tl.constexpr,
+    BLOCK_SIZE_N: tl.constexpr,
+):
+    # Map the program id to the elements of DW and DB it should compute.
+    pid = tl.program_id(0)
+    cols = pid * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
+    dw = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
+    db = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
+    # Iterate through the rows of DW and DB to sum the partial sums.
+    for i in range(0, M, BLOCK_SIZE_M):
+        rows = i + tl.arange(0, BLOCK_SIZE_M)
+        mask = (rows[:, None] < M) & (cols[None, :] < N)
+        offs = rows[:, None] * N + cols[None, :]
+        dw += tl.load(DW + offs, mask=mask, other=0.0)
+        db += tl.load(DB + offs, mask=mask, other=0.0)
+    # Write the final sum to the output.
+    sum_dw = tl.sum(dw, axis=0)
+    sum_db = tl.sum(db, axis=0)
+    tl.store(FINAL_DW + cols, sum_dw, mask=cols < N)
+    tl.store(FINAL_DB + cols, sum_db, mask=cols < N)
+```
+
+## Benchmark
+
+### Forward pass
+
+```bash
+python3 triton_layer_norm.py --mode forward
+layer-norm-forward:
+          N      Triton       Torch
+0    1024.0  528.516136  528.516136
+1    1536.0  572.158608  558.545450
+2    2048.0  595.781814  574.877182
+3    2560.0  602.352931  576.901391
+4    3072.0  592.192778  578.258826
+5    3584.0  616.602135  591.175279
+6    4096.0  624.524129  606.814814
+7    4608.0  609.322328  590.119034
+8    5120.0  615.939877  597.956224
+9    5632.0  621.462077  604.778525
+10   6144.0  630.153868  603.092009
+11   6656.0  630.153823  619.162791
+12   7168.0  633.635374  619.935145
+13   7680.0  639.999977  620.606079
+14   8192.0  642.509816  621.194295
+15   8704.0  641.769605  624.502255
+16   9216.0  643.912671  630.153829
+17   9728.0  643.173568  631.512128
+18  10240.0  645.039336  625.343543
+19  10752.0  637.155580  623.304334
+20  11264.0  641.366580  627.958183
+21  11776.0  643.058042  623.894039
+22  12288.0  644.616406  620.214515
+23  12800.0  644.025146  620.606042
+24  13312.0  645.430284  619.162791
+25  13824.0  646.736816  616.111424
+26  14336.0  647.954823  615.800311
+27  14848.0  652.547309  606.040836
+28  15360.0  650.158735  597.956224
+29  15872.0  651.158973  590.714847
+```
+![](./layer-norm-forward.png)
+
+### Backward pass
+
+```bash
+python3 triton_layer_norm.py --mode backward
+layer-norm-backward:
+          N      Triton       Torch
+0    1024.0  303.407414  351.085705
+1    1536.0  357.902918  400.695643
+2    2048.0  406.214885  413.042029
+3    2560.0  461.954908  393.846167
+4    3072.0  485.052653  359.648774
+5    3584.0  505.976473  353.975305
+6    4096.0  525.689846  351.085705
+7    4608.0  539.473160  341.333340
+8    5120.0  558.545450  331.212936
+9    5632.0  563.310028  329.678042
+10   6144.0  576.000001  317.793096
+11   6656.0  580.887276  321.416498
+12   7168.0  589.150662  319.762069
+13   7680.0  590.769251  319.445392
+14   8192.0  601.362644  314.070291
+15   8704.0  616.212406  309.934714
+16   9216.0  619.563005  305.068231
+17   9728.0  625.930315  302.817124
+18  10240.0  628.542186  299.342257
+19  10752.0  627.854035  297.290332
+20  11264.0  634.591530  297.726876
+21  11776.0  636.540534  295.014608
+22  12288.0  635.586191  291.703260
+23  12800.0  640.000016  291.737896
+24  13312.0  644.129041  291.769875
+25  13824.0  640.494226  293.347481
+26  14336.0  640.715067  291.834720
+27  14848.0  644.397807  288.544136
+28  15360.0  644.475536  288.225185
+29  15872.0  646.736839  287.059538
+```
+
+![](./layer-norm-backward.png)
+
+## Reference
+
+[Jimmy Lei Ba and Jamie Ryan Kiros and Geoffrey E. Hinton, "Layer Normalization", Arxiv 2016](https://arxiv.org/abs/1607.06450)
diff --git a/kernels/openai-triton/layer-norm/bwd-math.png b/kernels/openai-triton/layer-norm/bwd-math.png
new file mode 100644
index 0000000000000000000000000000000000000000..b2830d8d39609640e8189acffb74c06a76e07528
GIT binary patch
diff --git a/kernels/openai-triton/layer-norm/bwd.png b/kernels/openai-triton/layer-norm/bwd.png
new file mode 100644
index 0000000000000000000000000000000000000000..c531efc421c57f89152b24026ebbe9f09873de9a
GIT binary patch
diff --git a/kernels/openai-triton/layer-norm/layer-norm-backward.csv b/kernels/openai-triton/layer-norm/layer-norm-backward.csv
new file mode 100644
index 00000000..2207a6a3
--- /dev/null
+++ b/kernels/openai-triton/layer-norm/layer-norm-backward.csv
@@ -0,0 +1,31 @@
+N,Triton,Torch
+1024.000000,303.407414,351.085705
+1536.000000,357.902918,400.695643
+2048.000000,406.214885,413.042029
+2560.000000,461.954908,393.846167
+3072.000000,485.052653,359.648774
+3584.000000,505.976473,353.975305
+4096.000000,525.689846,351.085705
+4608.000000,539.473160,341.333340
+5120.000000,558.545450,331.212936
+5632.000000,563.310028,329.678042
+6144.000000,576.000001,317.793096
+6656.000000,580.887276,321.416498
+7168.000000,589.150662,319.762069
+7680.000000,590.769251,319.445392
+8192.000000,601.362644,314.070291
+8704.000000,616.212406,309.934714
+9216.000000,619.563005,305.068231
+9728.000000,625.930315,302.817124
+10240.000000,628.542186,299.342257
+10752.000000,627.854035,297.290332
+11264.000000,634.591530,297.726876
+11776.000000,636.540534,295.014608
+12288.000000,635.586191,291.703260
+12800.000000,640.000016,291.737896
+13312.000000,644.129041,291.769875
+13824.000000,640.494226,293.347481
+14336.000000,640.715067,291.834720
+14848.000000,644.397807,288.544136
+15360.000000,644.475536,288.225185
+15872.000000,646.736839,287.059538
diff --git a/kernels/openai-triton/layer-norm/layer-norm-backward.png b/kernels/openai-triton/layer-norm/layer-norm-backward.png
new file mode 100644
index 0000000000000000000000000000000000000000..4ee3cf7dd57dc11d348f43814e5254d7d0526e74
GIT binary patch
diff --git a/kernels/openai-triton/layer-norm/layer-norm-forward.csv b/kernels/openai-triton/layer-norm/layer-norm-forward.csv
new file mode 100644
index 00000000..d68ab5c8
--- /dev/null
+++ b/kernels/openai-triton/layer-norm/layer-norm-forward.csv
@@ -0,0 +1,31 @@
+N,Triton,Torch
+1024.000000,528.516136,528.516136
+1536.000000,572.158608,558.545450
+2048.000000,595.781814,574.877182
+2560.000000,602.352931,576.901391
+3072.000000,592.192778,578.258826
+3584.000000,616.602135,591.175279
+4096.000000,624.524129,606.814814
+4608.000000,609.322328,590.119034
+5120.000000,615.939877,597.956224
+5632.000000,621.462077,604.778525
+6144.000000,630.153868,603.092009
+6656.000000,630.153823,619.162791
+7168.000000,633.635374,619.935145
+7680.000000,639.999977,620.606079
+8192.000000,642.509816,621.194295
+8704.000000,641.769605,624.502255
+9216.000000,643.912671,630.153829
+9728.000000,643.173568,631.512128
+10240.000000,645.039336,625.343543
+10752.000000,637.155580,623.304334
+11264.000000,641.366580,627.958183
+11776.000000,643.058042,623.894039
+12288.000000,644.616406,620.214515
+12800.000000,644.025146,620.606042
+13312.000000,645.430284,619.162791
+13824.000000,646.736816,616.111424
+14336.000000,647.954823,615.800311
+14848.000000,652.547309,606.040836
+15360.000000,650.158735,597.956224
+15872.000000,651.158973,590.714847
diff --git a/kernels/openai-triton/layer-norm/layer-norm-forward.png b/kernels/openai-triton/layer-norm/layer-norm-forward.png
new file mode 100644
index 0000000000000000000000000000000000000000..6411243992ff3630e60b21f93945f8e994afb685
GIT binary patch
diff --git a/kernels/openai-triton/layer-norm/results.html b/kernels/openai-triton/layer-norm/results.html
new file mode 100644
index 00000000..e49c4b28
--- /dev/null
+++ b/kernels/openai-triton/layer-norm/results.html
@@ -0,0 +1,3 @@
+<html><body>
+<image src="layer-norm-backward.png"/>
+</body></html>
diff --git a/kernels/openai-triton/layer-norm/triton_layer_norm.py b/kernels/openai-triton/layer-norm/triton_layer_norm.py
new file mode 100644
index 00000000..375d5dda
--- /dev/null
+++ b/kernels/openai-triton/layer-norm/triton_layer_norm.py
@@ -0,0 +1,401 @@
+"""
+Layer Normalization
+====================
+In this tutorial, you will write a high-performance layer normalization
+kernel that runs faster than the PyTorch implementation.
+
+In doing so, you will learn about:
+
+* Implementing backward pass in Triton.
+
+* Implementing parallel reduction in Triton.
+
+"""
+
+# Adapted from: https://github.com/triton-lang/triton/tree/main/python/turorials/05-layer-norm.py
+
+import argparse
+
+import torch
+import triton
+import triton.language as tl
+
+try:
+    # This is https://github.com/NVIDIA/apex, NOT the apex on PyPi, so it
+    # should not be added to extras_require in setup.py.
+    import apex
+
+    HAS_APEX = True
+except ModuleNotFoundError:
+    HAS_APEX = False
+
+DEVICE = triton.runtime.driver.active.get_active_torch_device()
+
+
+@triton.jit
+def _layer_norm_fwd_fused(
+    X,  # pointer to the input
+    Y,  # pointer to the output
+    W,  # pointer to the weights
+    B,  # pointer to the biases
+    Mean,  # pointer to the mean
+    Rstd,  # pointer to the 1/std
+    stride,  # how much to increase the pointer when moving by 1 row
+    N,  # number of columns in X
+    eps,  # epsilon to avoid division by zero
+    BLOCK_SIZE: tl.constexpr,
+):
+    # Map the program id to the row of X and Y it should compute.
+    row = tl.program_id(0)
+    Y += row * stride
+    X += row * stride
+    # Compute mean
+    mean = 0
+    _mean = tl.zeros([BLOCK_SIZE], dtype=tl.float32)
+    for off in range(0, N, BLOCK_SIZE):
+        cols = off + tl.arange(0, BLOCK_SIZE)
+        a = tl.load(X + cols, mask=cols < N, other=0.0).to(tl.float32)
+        _mean += a
+    mean = tl.sum(_mean, axis=0) / N
+    # Compute variance
+    _var = tl.zeros([BLOCK_SIZE], dtype=tl.float32)
+    for off in range(0, N, BLOCK_SIZE):
+        cols = off + tl.arange(0, BLOCK_SIZE)
+        x = tl.load(X + cols, mask=cols < N, other=0.0).to(tl.float32)
+        x = tl.where(cols < N, x - mean, 0.0)
+        _var += x * x
+    var = tl.sum(_var, axis=0) / N
+    rstd = 1 / tl.sqrt(var + eps)
+    # Write mean / rstd
+    tl.store(Mean + row, mean)
+    tl.store(Rstd + row, rstd)
+    # Normalize and apply linear transformation
+    for off in range(0, N, BLOCK_SIZE):
+        cols = off + tl.arange(0, BLOCK_SIZE)
+        mask = cols < N
+        w = tl.load(W + cols, mask=mask)
+        b = tl.load(B + cols, mask=mask)
+        x = tl.load(X + cols, mask=mask, other=0.0).to(tl.float32)
+        x_hat = (x - mean) * rstd
+        y = x_hat * w + b
+        # Write output
+        tl.store(Y + cols, y, mask=mask)
+
+
+@triton.jit
+def _layer_norm_bwd_dx_fused(
+    DX,  # pointer to the input gradient
+    DY,  # pointer to the output gradient
+    DW,  # pointer to the partial sum of weights gradient
+    DB,  # pointer to the partial sum of biases gradient
+    X,  # pointer to the input
+    W,  # pointer to the weights
+    Mean,  # pointer to the mean
+    Rstd,  # pointer to the 1/std
+    Lock,  # pointer to the lock
+    stride,  # how much to increase the pointer when moving by 1 row
+    N,  # number of columns in X
+    GROUP_SIZE_M: tl.constexpr,
+    BLOCK_SIZE_N: tl.constexpr,
+):
+    # Map the program id to the elements of X, DX, and DY it should compute.
+    row = tl.program_id(0)
+    cols = tl.arange(0, BLOCK_SIZE_N)
+    mask = cols < N
+    X += row * stride
+    DY += row * stride
+    DX += row * stride
+    # Offset locks and weights/biases gradient pointer for parallel reduction
+    lock_id = row % GROUP_SIZE_M
+    Lock += lock_id
+    Count = Lock + GROUP_SIZE_M
+    DW = DW + lock_id * N + cols
+    DB = DB + lock_id * N + cols
+    # Load data to SRAM
+    x = tl.load(X + cols, mask=mask, other=0).to(tl.float32)
+    dy = tl.load(DY + cols, mask=mask, other=0).to(tl.float32)
+    w = tl.load(W + cols, mask=mask).to(tl.float32)
+    mean = tl.load(Mean + row)
+    rstd = tl.load(Rstd + row)
+    # Compute dx
+    xhat = (x - mean) * rstd
+    wdy = w * dy
+    xhat = tl.where(mask, xhat, 0.0)
+    wdy = tl.where(mask, wdy, 0.0)
+    c1 = tl.sum(xhat * wdy, axis=0) / N
+    c2 = tl.sum(wdy, axis=0) / N
+    dx = (wdy - (xhat * c1 + c2)) * rstd
+    # Write dx
+    tl.store(DX + cols, dx, mask=mask)
+    # Accumulate partial sums for dw/db
+    partial_dw = (dy * xhat).to(w.dtype)
+    partial_db = (dy).to(w.dtype)
+    while tl.atomic_cas(Lock, 0, 1) == 1:
+        pass
+    count = tl.load(Count)
+    # First store doesn't accumulate
+    if count == 0:
+        tl.atomic_xchg(Count, 1)
+    else:
+        partial_dw += tl.load(DW, mask=mask)
+        partial_db += tl.load(DB, mask=mask)
+    tl.store(DW, partial_dw, mask=mask)
+    tl.store(DB, partial_db, mask=mask)
+
+    # need a barrier to ensure all threads finished before
+    # releasing the lock
+    tl.debug_barrier()
+
+    # Release the lock
+    tl.atomic_xchg(Lock, 0)
+
+
+@triton.jit
+def _layer_norm_bwd_dwdb(
+    DW,  # pointer to the partial sum of weights gradient
+    DB,  # pointer to the partial sum of biases gradient
+    FINAL_DW,  # pointer to the weights gradient
+    FINAL_DB,  # pointer to the biases gradient
+    M,  # GROUP_SIZE_M
+    N,  # number of columns
+    BLOCK_SIZE_M: tl.constexpr,
+    BLOCK_SIZE_N: tl.constexpr,
+):
+    # Map the program id to the elements of DW and DB it should compute.
+    pid = tl.program_id(0)
+    cols = pid * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
+    dw = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
+    db = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
+    # Iterate through the rows of DW and DB to sum the partial sums.
+    for i in range(0, M, BLOCK_SIZE_M):
+        rows = i + tl.arange(0, BLOCK_SIZE_M)
+        mask = (rows[:, None] < M) & (cols[None, :] < N)
+        offs = rows[:, None] * N + cols[None, :]
+        dw += tl.load(DW + offs, mask=mask, other=0.0)
+        db += tl.load(DB + offs, mask=mask, other=0.0)
+    # Write the final sum to the output.
+    sum_dw = tl.sum(dw, axis=0)
+    sum_db = tl.sum(db, axis=0)
+    tl.store(FINAL_DW + cols, sum_dw, mask=cols < N)
+    tl.store(FINAL_DB + cols, sum_db, mask=cols < N)
+
+
+class LayerNorm(torch.autograd.Function):
+
+    @staticmethod
+    def forward(ctx, x, normalized_shape, weight, bias, eps):
+        # allocate output
+        y = torch.empty_like(x)
+        # reshape input data into 2D tensor
+        x_arg = x.reshape(-1, x.shape[-1])
+        M, N = x_arg.shape
+        mean = torch.empty((M,), dtype=torch.float32, device=x.device)
+        rstd = torch.empty((M,), dtype=torch.float32, device=x.device)
+        # Less than 64KB per feature: enqueue fused kernel
+        MAX_FUSED_SIZE = 65536 // x.element_size()
+        BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(N))
+        if N > BLOCK_SIZE:
+            raise RuntimeError(
+                "This layer norm doesn't support feature dim >= 64KB."
+            )
+        # heuristics for number of warps
+        num_warps = min(max(BLOCK_SIZE // 256, 1), 8)
+        # enqueue kernel
+        _layer_norm_fwd_fused[(M,)](  #
+            x_arg,
+            y,
+            weight,
+            bias,
+            mean,
+            rstd,  #
+            x_arg.stride(0),
+            N,
+            eps,  #
+            BLOCK_SIZE=BLOCK_SIZE,
+            num_warps=num_warps,
+            num_ctas=1,
+        )
+        ctx.save_for_backward(x, weight, bias, mean, rstd)
+        ctx.BLOCK_SIZE = BLOCK_SIZE
+        ctx.num_warps = num_warps
+        ctx.eps = eps
+        return y
+
+    @staticmethod
+    def backward(ctx, dy):
+        x, w, b, m, v = ctx.saved_tensors
+        # heuristics for amount of parallel reduction stream for DW/DB
+        N = w.shape[0]
+        GROUP_SIZE_M = 64
+        if N <= 8192:
+            GROUP_SIZE_M = 96
+        if N <= 4096:
+            GROUP_SIZE_M = 128
+        if N <= 1024:
+            GROUP_SIZE_M = 256
+        # allocate output
+        locks = torch.zeros(
+            2 * GROUP_SIZE_M, dtype=torch.int32, device=w.device
+        )
+        _dw = torch.zeros((GROUP_SIZE_M, N), dtype=x.dtype, device=w.device)
+        _db = torch.zeros((GROUP_SIZE_M, N), dtype=x.dtype, device=w.device)
+        dw = torch.empty((N,), dtype=w.dtype, device=w.device)
+        db = torch.empty((N,), dtype=w.dtype, device=w.device)
+        dx = torch.empty_like(dy)
+        # enqueue kernel using forward pass heuristics
+        # also compute partial sums for DW and DB
+        x_arg = x.reshape(-1, x.shape[-1])
+        M, N = x_arg.shape
+        _layer_norm_bwd_dx_fused[(M,)](  #
+            dx,
+            dy,
+            _dw,
+            _db,
+            x,
+            w,
+            m,
+            v,
+            locks,  #
+            x_arg.stride(0),
+            N,  #
+            BLOCK_SIZE_N=ctx.BLOCK_SIZE,  #
+            GROUP_SIZE_M=GROUP_SIZE_M,  #
+            num_warps=ctx.num_warps,
+        )
+        grid = lambda meta: (triton.cdiv(N, meta["BLOCK_SIZE_N"]),)
+        # accumulate partial sums in separate kernel
+        _layer_norm_bwd_dwdb[grid](
+            _dw,
+            _db,
+            dw,
+            db,
+            min(GROUP_SIZE_M, M),
+            N,  #
+            BLOCK_SIZE_M=32,  #
+            BLOCK_SIZE_N=128,
+            num_ctas=1,
+        )
+        return dx, None, dw, db, None
+
+
+layer_norm = LayerNorm.apply
+
+
+def test_layer_norm(M, N, dtype, eps=1e-5, device=DEVICE):
+    # create data
+    x_shape = (M, N)
+    w_shape = (x_shape[-1],)
+    weight = torch.rand(w_shape, dtype=dtype, device=device, requires_grad=True)
+    bias = torch.rand(w_shape, dtype=dtype, device=device, requires_grad=True)
+    x = -2.3 + 0.5 * torch.randn(x_shape, dtype=dtype, device=device)
+    dy = 0.1 * torch.randn_like(x)
+    x.requires_grad_(True)
+    # forward pass
+    y_tri = layer_norm(x, w_shape, weight, bias, eps)
+    y_ref = torch.nn.functional.layer_norm(x, w_shape, weight, bias, eps).to(
+        dtype
+    )
+    # backward pass (triton)
+    y_tri.backward(dy, retain_graph=True)
+    dx_tri, dw_tri, db_tri = [_.grad.clone() for _ in [x, weight, bias]]
+    x.grad, weight.grad, bias.grad = None, None, None
+    # backward pass (torch)
+    y_ref.backward(dy, retain_graph=True)
+    dx_ref, dw_ref, db_ref = [_.grad.clone() for _ in [x, weight, bias]]
+    # compare
+    assert torch.allclose(y_tri, y_ref, atol=1e-2, rtol=0)
+    assert torch.allclose(dx_tri, dx_ref, atol=1e-2, rtol=0)
+    assert torch.allclose(db_tri, db_ref, atol=1e-2, rtol=0)
+    assert torch.allclose(dw_tri, dw_ref, atol=1e-2, rtol=0)
+
+
+def get_args():
+    parser = argparse.ArgumentParser(
+        description="Run layer norm benchmark",
+        formatter_class=argparse.ArgumentDefaultsHelpFormatter,
+    )
+    parser.add_argument(
+        "--mode",
+        type=str,
+        default="forward",
+        choices=["forward", "backward"],
+        help="Benchmark mode",
+    )
+    return parser.parse_args()
+
+
+args = get_args()
+if args.mode not in ["forward", "backward"]:
+    raise ValueError(
+        f"Invalid mode: {args.mode}. Choose 'forward' or 'backward'."
+    )
+bench_mode = args.mode
+
+
+@triton.testing.perf_report(
+    triton.testing.Benchmark(
+        x_names=["N"],
+        x_vals=[512 * i for i in range(2, 32)],
+        line_arg="provider",
+        line_vals=["triton", "torch"] + (["apex"] if HAS_APEX else []),
+        line_names=["Triton", "Torch"] + (["Apex"] if HAS_APEX else []),
+        styles=[("blue", "-"), ("green", "-"), ("orange", "-")],
+        ylabel="GB/s",
+        plot_name=f"layer-norm-{bench_mode}",
+        args={"M": 4096, "dtype": torch.float16, "mode": bench_mode},
+    )
+)
+def bench_layer_norm(
+    M, N, dtype, provider, mode="forward", eps=1e-5, device=DEVICE
+):
+    # create data
+    x_shape = (M, N)
+    w_shape = (x_shape[-1],)
+    weight = torch.rand(w_shape, dtype=dtype, device=device, requires_grad=True)
+    bias = torch.rand(w_shape, dtype=dtype, device=device, requires_grad=True)
+    x = -2.3 + 0.5 * torch.randn(x_shape, dtype=dtype, device=device)
+    dy = 0.1 * torch.randn_like(x)
+    x.requires_grad_(True)
+    quantiles = [0.5, 0.2, 0.8]
+
+    def y_fwd():
+
+        if provider == "triton":
+            return layer_norm(x, w_shape, weight, bias, eps)  # noqa: F811, E704
+
+        if provider == "torch":
+            return torch.nn.functional.layer_norm(
+                x, w_shape, weight, bias, eps
+            )  # noqa: F811, E704
+
+        if provider == "apex":
+            apex_layer_norm = (
+                apex.normalization.FusedLayerNorm(w_shape)
+                .to(x.device)
+                .to(x.dtype)
+            )
+            return apex_layer_norm(x)  # noqa: F811, E704
+
+    # forward pass
+    if mode == "forward":
+        gbps = lambda ms: 2 * x.numel() * x.element_size() * 1e-9 / (ms * 1e-3)
+        ms, min_ms, max_ms = triton.testing.do_bench(
+            y_fwd, quantiles=quantiles, rep=500
+        )
+    # backward pass
+    if mode == "backward":
+        y = y_fwd()
+        gbps = (
+            lambda ms: 3 * x.numel() * x.element_size() * 1e-9 / (ms * 1e-3)
+        )  # noqa: F811, E704
+        ms, min_ms, max_ms = triton.testing.do_bench(
+            lambda: y.backward(dy, retain_graph=True),
+            quantiles=quantiles,
+            grad_to_none=[x],
+            rep=500,
+        )
+    return gbps(ms), gbps(max_ms), gbps(min_ms)
+
+
+test_layer_norm(1151, 8192, torch.float16)
+bench_layer_norm.run(save_path=".", print_data=True)
diff --git a/kernels/openai-triton/flash-mla/.gitignore b/kernels/openai-triton/matrix-multiplication/.gitignore
similarity index 100%
rename from kernels/openai-triton/flash-mla/.gitignore
rename to kernels/openai-triton/matrix-multiplication/.gitignore
diff --git a/kernels/openai-triton/reduce/.gitignore b/kernels/openai-triton/reduce/.gitignore
deleted file mode 100644
index a9886a40..00000000
--- a/kernels/openai-triton/reduce/.gitignore
+++ /dev/null
@@ -1,20 +0,0 @@
-*.so
-*.a
-*.dylib
-*.dll
-*.lib
-.DS_Store
-build
-*.whl
-tmp
-*.onnx
-*.pb
-*.pbtxt
-*.pt
-*.pth
-*.engine
-*.bin
-*.nsys
-*.nvvp
-*.nsys*
-*.sqlite
diff --git a/kernels/openai-triton/sgemm/.gitignore b/kernels/openai-triton/sgemm/.gitignore
deleted file mode 100644
index a9886a40..00000000
--- a/kernels/openai-triton/sgemm/.gitignore
+++ /dev/null
@@ -1,20 +0,0 @@
-*.so
-*.a
-*.dylib
-*.dll
-*.lib
-.DS_Store
-build
-*.whl
-tmp
-*.onnx
-*.pb
-*.pbtxt
-*.pt
-*.pth
-*.engine
-*.bin
-*.nsys
-*.nvvp
-*.nsys*
-*.sqlite
diff --git a/kernels/openai-triton/elementwise/.gitignore b/kernels/openai-triton/vector-add/.gitignore
similarity index 100%
rename from kernels/openai-triton/elementwise/.gitignore
rename to kernels/openai-triton/vector-add/.gitignore
diff --git a/kernels/openai-triton/elementwise/README.md b/kernels/openai-triton/vector-add/README.md
similarity index 100%
rename from kernels/openai-triton/elementwise/README.md
rename to kernels/openai-triton/vector-add/README.md
diff --git a/kernels/openai-triton/elementwise/results.html b/kernels/openai-triton/vector-add/results.html
similarity index 100%
rename from kernels/openai-triton/elementwise/results.html
rename to kernels/openai-triton/vector-add/results.html
diff --git a/kernels/openai-triton/elementwise/triton_vector_add.py b/kernels/openai-triton/vector-add/triton_vector_add.py
similarity index 100%
rename from kernels/openai-triton/elementwise/triton_vector_add.py
rename to kernels/openai-triton/vector-add/triton_vector_add.py
diff --git a/kernels/openai-triton/elementwise/vector-add-performance.png b/kernels/openai-triton/vector-add/vector-add-performance.png
similarity index 100%
rename from kernels/openai-triton/elementwise/vector-add-performance.png
rename to kernels/openai-triton/vector-add/vector-add-performance.png

From 7289e6537c783c01f991bf61453f63537aad7e1a Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Sun, 29 Jun 2025 10:04:44 +0800
Subject: [PATCH 38/57] fix flash-attn comments (#354)

* Update README.md

* Update flash_attn_mma_share_kv.cu

* fix flash-attn comments
---
 README.md                                                 | 6 ++----
 kernels/flash-attn/mma/basic/flash_attn_mma_share_kv.cu   | 8 ++++----
 .../mma/basic/flash_attn_mma_share_kv_F32F16F16F32.cu     | 8 ++++----
 kernels/flash-attn/mma/basic/flash_attn_mma_share_qkv.cu  | 8 ++++----
 .../mma/basic/flash_attn_mma_share_qkv_F32F16F16F32.cu    | 8 ++++----
 kernels/flash-attn/mma/basic/flash_attn_mma_split_kv.cu   | 8 ++++----
 kernels/flash-attn/mma/basic/flash_attn_mma_split_q.cu    | 8 ++++----
 kernels/flash-attn/mma/basic/flash_attn_mma_tiling_qk.cu  | 8 ++++----
 .../mma/basic/flash_attn_mma_tiling_qk_F32F16F16F32.cu    | 8 ++++----
 kernels/flash-attn/mma/basic/flash_attn_mma_tiling_qkv.cu | 8 ++++----
 .../mma/basic/flash_attn_mma_tiling_qkv_F32F16F16F32.cu   | 8 ++++----
 .../mma/others/flash_attn_mma_share_kv_F32F16F16F32_rr.cu | 8 ++++----
 .../others/flash_attn_mma_share_qkv_F32F16F16F32_rr.cu    | 8 ++++----
 .../mma/others/flash_attn_mma_share_qkv_Os2g.cu           | 8 ++++----
 .../others/flash_attn_mma_tiling_qk_F32F16F16F32_rr.cu    | 8 ++++----
 .../mma/swizzle/flash_attn_mma_share_kv_swizzle_q.cu      | 8 ++++----
 .../mma/swizzle/flash_attn_mma_share_kv_swizzle_qk.cu     | 8 ++++----
 .../mma/swizzle/flash_attn_mma_share_kv_swizzle_qkv.cu    | 8 ++++----
 .../mma/swizzle/flash_attn_mma_share_qkv_swizzle_q.cu     | 8 ++++----
 .../mma/swizzle/flash_attn_mma_share_qkv_swizzle_qk.cu    | 8 ++++----
 .../mma/swizzle/flash_attn_mma_share_qkv_swizzle_qkv.cu   | 8 ++++----
 .../mma/swizzle/flash_attn_mma_tiling_qk_swizzle_q.cu     | 8 ++++----
 .../mma/swizzle/flash_attn_mma_tiling_qk_swizzle_qk.cu    | 8 ++++----
 .../mma/swizzle/flash_attn_mma_tiling_qk_swizzle_qkv.cu   | 8 ++++----
 .../mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_q.cu    | 8 ++++----
 .../flash_attn_mma_tiling_qkv_swizzle_q_F32F16F16F32.cu   | 8 ++++----
 .../mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_qk.cu   | 8 ++++----
 .../flash_attn_mma_tiling_qkv_swizzle_qk_F32F16F16F32.cu  | 8 ++++----
 .../mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_qkv.cu  | 8 ++++----
 .../flash_attn_mma_tiling_qkv_swizzle_qkv_F32F16F16F32.cu | 8 ++++----
 30 files changed, 118 insertions(+), 120 deletions(-)

diff --git a/README.md b/README.md
index c7451416..75bbd459 100644
--- a/README.md
+++ b/README.md
@@ -7,10 +7,8 @@
   <div align='center'>
       <img src=https://cdn.rawgit.com/sindresorhus/awesome/d7305f38d29fed78fa85652e3a63e154dd8e8829/media/badge.svg >
       <img src=https://img.shields.io/badge/Language-CUDA-brightgreen.svg >
-      <img src=https://img.shields.io/github/watchers/xlite-dev/LeetCUDA?color=9cc >
-      <img src=https://img.shields.io/github/forks/xlite-dev/LeetCUDA.svg?style=social >
-      <img src=https://img.shields.io/github/stars/xlite-dev/LeetCUDA.svg?style=social >
-      <img src=https://img.shields.io/badge/Release-v3.0.12-brightgreen.svg >
+      <img src=https://img.shields.io/github/forks/xlite-dev/LeetCUDA.svg?style=dark >
+      <img src=https://img.shields.io/github/stars/xlite-dev/LeetCUDA.svg?style=dark >
       <img src=https://img.shields.io/badge/License-GPLv3.0-turquoise.svg >
   </div>
 </div>
diff --git a/kernels/flash-attn/mma/basic/flash_attn_mma_share_kv.cu b/kernels/flash-attn/mma/basic/flash_attn_mma_share_kv.cu
index 9981ce0a..4a5d8dc7 100644
--- a/kernels/flash-attn/mma/basic/flash_attn_mma_share_kv.cu
+++ b/kernels/flash-attn/mma/basic/flash_attn_mma_share_kv.cu
@@ -173,16 +173,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d]. Allocate R_Q[(kHeadDim/kMmaAtomK)<=8][1][4],
   // e.g R_Q[4][1][4] 16 regs. By the way, we have to reduce R_Z to 0 regs and
   // reuse R_Q for collective store. Then we can load Q from smem only once and
diff --git a/kernels/flash-attn/mma/basic/flash_attn_mma_share_kv_F32F16F16F32.cu b/kernels/flash-attn/mma/basic/flash_attn_mma_share_kv_F32F16F16F32.cu
index 4fdc0bee..c6e17559 100644
--- a/kernels/flash-attn/mma/basic/flash_attn_mma_share_kv_F32F16F16F32.cu
+++ b/kernels/flash-attn/mma/basic/flash_attn_mma_share_kv_F32F16F16F32.cu
@@ -174,16 +174,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d]. Allocate R_Q[(kHeadDim/kMmaAtomK)<=8][1][4],
   // e.g R_Q[4][1][4] 16 regs. By the way, we have to reduce R_Z to 0 regs and
   // reuse R_Q for collective store. Then we can load Q from smem only once and
diff --git a/kernels/flash-attn/mma/basic/flash_attn_mma_share_qkv.cu b/kernels/flash-attn/mma/basic/flash_attn_mma_share_qkv.cu
index 92de2a6d..0a8b12f2 100644
--- a/kernels/flash-attn/mma/basic/flash_attn_mma_share_qkv.cu
+++ b/kernels/flash-attn/mma/basic/flash_attn_mma_share_qkv.cu
@@ -166,16 +166,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d]. Allocate R_Q[(kHeadDim/kMmaAtomK)<=8][1][4],
   // e.g R_Q[4][1][4] 16 regs. By the way, we have to reduce R_Z to 0 regs and
   // reuse R_Q for collective store. Then we can load Q from smem only once and
diff --git a/kernels/flash-attn/mma/basic/flash_attn_mma_share_qkv_F32F16F16F32.cu b/kernels/flash-attn/mma/basic/flash_attn_mma_share_qkv_F32F16F16F32.cu
index 655e3c7d..be9388bf 100644
--- a/kernels/flash-attn/mma/basic/flash_attn_mma_share_qkv_F32F16F16F32.cu
+++ b/kernels/flash-attn/mma/basic/flash_attn_mma_share_qkv_F32F16F16F32.cu
@@ -167,16 +167,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d]. Allocate R_Q[(kHeadDim/kMmaAtomK)<=8][1][4],
   // e.g R_Q[4][1][4] 16 regs. By the way, we have to reduce R_Z to 0 regs and
   // reuse R_Q for collective store. Then we can load Q from smem only once and
diff --git a/kernels/flash-attn/mma/basic/flash_attn_mma_split_kv.cu b/kernels/flash-attn/mma/basic/flash_attn_mma_split_kv.cu
index e2c9aa31..8f08b3cf 100644
--- a/kernels/flash-attn/mma/basic/flash_attn_mma_split_kv.cu
+++ b/kernels/flash-attn/mma/basic/flash_attn_mma_split_kv.cu
@@ -133,8 +133,8 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
   uint32_t smem_S_base_ptr = __cvta_generic_to_shared(S_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2];
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2];
@@ -146,8 +146,8 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   __shared__ float block_row_max_new_smem[Br][kMmaTileSeqLenK + 1];
   __shared__ float block_row_sum_new_smem[Br][kMmaTileSeqLenK + 1];
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d].
   uint32_t R_Q[kWarpTileSeqLenQ][4];
   uint32_t R_K[kWarpTileSeqLenK][2];
diff --git a/kernels/flash-attn/mma/basic/flash_attn_mma_split_q.cu b/kernels/flash-attn/mma/basic/flash_attn_mma_split_q.cu
index 8b1860c8..9ad7d5b6 100644
--- a/kernels/flash-attn/mma/basic/flash_attn_mma_split_q.cu
+++ b/kernels/flash-attn/mma/basic/flash_attn_mma_split_q.cu
@@ -154,16 +154,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d].
   uint32_t R_Q[kWarpTileSeqLenQ][4];  // [1][4]
   uint32_t R_K[kWarpTileSeqLenK][2];  // [8][2]
diff --git a/kernels/flash-attn/mma/basic/flash_attn_mma_tiling_qk.cu b/kernels/flash-attn/mma/basic/flash_attn_mma_tiling_qk.cu
index c32e2f1b..87e39db3 100644
--- a/kernels/flash-attn/mma/basic/flash_attn_mma_tiling_qk.cu
+++ b/kernels/flash-attn/mma/basic/flash_attn_mma_tiling_qk.cu
@@ -187,16 +187,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d].
   uint32_t R_Q[kWarpTileSeqLenQ][4];  // [1][4]
   uint32_t R_K[kWarpTileSeqLenK][2];  // [8][2]
diff --git a/kernels/flash-attn/mma/basic/flash_attn_mma_tiling_qk_F32F16F16F32.cu b/kernels/flash-attn/mma/basic/flash_attn_mma_tiling_qk_F32F16F16F32.cu
index 0d63977f..6cdaebe6 100644
--- a/kernels/flash-attn/mma/basic/flash_attn_mma_tiling_qk_F32F16F16F32.cu
+++ b/kernels/flash-attn/mma/basic/flash_attn_mma_tiling_qk_F32F16F16F32.cu
@@ -188,16 +188,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d].
   uint32_t R_Q[kWarpTileSeqLenQ][4];  // [1][4]
   uint32_t R_K[kWarpTileSeqLenK][2];  // [8][2]
diff --git a/kernels/flash-attn/mma/basic/flash_attn_mma_tiling_qkv.cu b/kernels/flash-attn/mma/basic/flash_attn_mma_tiling_qkv.cu
index 31feee48..935c1d6d 100644
--- a/kernels/flash-attn/mma/basic/flash_attn_mma_tiling_qkv.cu
+++ b/kernels/flash-attn/mma/basic/flash_attn_mma_tiling_qkv.cu
@@ -170,16 +170,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d].
   uint32_t R_Q[kWarpTileSeqLenQ][4]; // [1][4]
   uint32_t R_K[kWarpTileSeqLenK][2]; // [8][2]
diff --git a/kernels/flash-attn/mma/basic/flash_attn_mma_tiling_qkv_F32F16F16F32.cu b/kernels/flash-attn/mma/basic/flash_attn_mma_tiling_qkv_F32F16F16F32.cu
index 06dab5b6..42892f18 100644
--- a/kernels/flash-attn/mma/basic/flash_attn_mma_tiling_qkv_F32F16F16F32.cu
+++ b/kernels/flash-attn/mma/basic/flash_attn_mma_tiling_qkv_F32F16F16F32.cu
@@ -171,16 +171,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d].
   uint32_t R_Q[kWarpTileSeqLenQ][4]; // [1][4]
   uint32_t R_K[kWarpTileSeqLenK][2]; // [8][2]
diff --git a/kernels/flash-attn/mma/others/flash_attn_mma_share_kv_F32F16F16F32_rr.cu b/kernels/flash-attn/mma/others/flash_attn_mma_share_kv_F32F16F16F32_rr.cu
index 6f934315..0cfd2d66 100644
--- a/kernels/flash-attn/mma/others/flash_attn_mma_share_kv_F32F16F16F32_rr.cu
+++ b/kernels/flash-attn/mma/others/flash_attn_mma_share_kv_F32F16F16F32_rr.cu
@@ -174,16 +174,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d]. Allocate R_Q[(kHeadDim/kMmaAtomK)<=8][1][4],
   // e.g R_Q[4][1][4] 16 regs. By the way, we have to reduce R_Z to 0 regs and
   // reuse R_Q for collective store. Then we can load Q from smem only once and
diff --git a/kernels/flash-attn/mma/others/flash_attn_mma_share_qkv_F32F16F16F32_rr.cu b/kernels/flash-attn/mma/others/flash_attn_mma_share_qkv_F32F16F16F32_rr.cu
index 547433bb..cd7e5d0f 100644
--- a/kernels/flash-attn/mma/others/flash_attn_mma_share_qkv_F32F16F16F32_rr.cu
+++ b/kernels/flash-attn/mma/others/flash_attn_mma_share_qkv_F32F16F16F32_rr.cu
@@ -168,16 +168,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d]. Allocate R_Q[(kHeadDim/kMmaAtomK)<=8][1][4],
   // e.g R_Q[4][1][4] 16 regs. By the way, we have to reduce R_Z to 0 regs and
   // reuse R_Q for collective store. Then we can load Q from smem only once and
diff --git a/kernels/flash-attn/mma/others/flash_attn_mma_share_qkv_Os2g.cu b/kernels/flash-attn/mma/others/flash_attn_mma_share_qkv_Os2g.cu
index c3c99804..844363e0 100644
--- a/kernels/flash-attn/mma/others/flash_attn_mma_share_qkv_Os2g.cu
+++ b/kernels/flash-attn/mma/others/flash_attn_mma_share_qkv_Os2g.cu
@@ -162,16 +162,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d]. Allocate R_Q[(kHeadDim/kMmaAtomK)<=8][1][4],
   // e.g R_Q[4][1][4] 16 regs. By the way, we have to reduce R_Z to 0 regs and
   // reuse R_Q for collective store. Then we can load Q from smem only once and
diff --git a/kernels/flash-attn/mma/others/flash_attn_mma_tiling_qk_F32F16F16F32_rr.cu b/kernels/flash-attn/mma/others/flash_attn_mma_tiling_qk_F32F16F16F32_rr.cu
index ad2b4da1..3fe6246f 100644
--- a/kernels/flash-attn/mma/others/flash_attn_mma_tiling_qk_F32F16F16F32_rr.cu
+++ b/kernels/flash-attn/mma/others/flash_attn_mma_tiling_qk_F32F16F16F32_rr.cu
@@ -172,16 +172,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d].
   uint32_t R_Q[kWarpTileSeqLenQ][4]; // [1][4]
   uint32_t R_K[kWarpTileSeqLenK][2]; // [8][2]
diff --git a/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_kv_swizzle_q.cu b/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_kv_swizzle_q.cu
index 12124286..7f8493f5 100644
--- a/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_kv_swizzle_q.cu
+++ b/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_kv_swizzle_q.cu
@@ -227,16 +227,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d]. Allocate R_Q[(kHeadDim/kMmaAtomK)<=8][1][4],
   // e.g R_Q[4][1][4] 16 regs. By the way, we have to reduce R_Z to 0 regs and
   // reuse R_Q for collective store. Then we can load Q from smem only once and
diff --git a/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_kv_swizzle_qk.cu b/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_kv_swizzle_qk.cu
index f79df865..101524ff 100644
--- a/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_kv_swizzle_qk.cu
+++ b/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_kv_swizzle_qk.cu
@@ -228,16 +228,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d]. Allocate R_Q[(kHeadDim/kMmaAtomK)<=8][1][4],
   // e.g R_Q[4][1][4] 16 regs. By the way, we have to reduce R_Z to 0 regs and
   // reuse R_Q for collective store. Then we can load Q from smem only once and
diff --git a/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_kv_swizzle_qkv.cu b/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_kv_swizzle_qkv.cu
index 334579c7..254a9645 100644
--- a/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_kv_swizzle_qkv.cu
+++ b/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_kv_swizzle_qkv.cu
@@ -231,16 +231,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d]. Allocate R_Q[(kHeadDim/kMmaAtomK)<=8][1][4],
   // e.g R_Q[4][1][4] 16 regs. By the way, we have to reduce R_Z to 0 regs and
   // reuse R_Q for collective store. Then we can load Q from smem only once and
diff --git a/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_qkv_swizzle_q.cu b/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_qkv_swizzle_q.cu
index fcf98e7b..aefd62f8 100644
--- a/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_qkv_swizzle_q.cu
+++ b/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_qkv_swizzle_q.cu
@@ -225,16 +225,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d]. Allocate R_Q[(kHeadDim/kMmaAtomK)<=8][1][4],
   // e.g R_Q[4][1][4] 16 regs. By the way, we have to reduce R_Z to 0 regs and
   // reuse R_Q for collective store. Then we can load Q from smem only once and
diff --git a/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_qkv_swizzle_qk.cu b/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_qkv_swizzle_qk.cu
index ef06340f..5dd1194c 100644
--- a/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_qkv_swizzle_qk.cu
+++ b/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_qkv_swizzle_qk.cu
@@ -226,16 +226,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d]. Allocate R_Q[(kHeadDim/kMmaAtomK)<=8][1][4],
   // e.g R_Q[4][1][4] 16 regs. By the way, we have to reduce R_Z to 0 regs and
   // reuse R_Q for collective store. Then we can load Q from smem only once and
diff --git a/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_qkv_swizzle_qkv.cu b/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_qkv_swizzle_qkv.cu
index 675889af..2f424ede 100644
--- a/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_qkv_swizzle_qkv.cu
+++ b/kernels/flash-attn/mma/swizzle/flash_attn_mma_share_qkv_swizzle_qkv.cu
@@ -227,16 +227,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d]. Allocate R_Q[(kHeadDim/kMmaAtomK)<=8][1][4],
   // e.g R_Q[4][1][4] 16 regs. By the way, we have to reduce R_Z to 0 regs and
   // reuse R_Q for collective store. Then we can load Q from smem only once and
diff --git a/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qk_swizzle_q.cu b/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qk_swizzle_q.cu
index 241c929f..eff14167 100644
--- a/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qk_swizzle_q.cu
+++ b/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qk_swizzle_q.cu
@@ -324,16 +324,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d].
   uint32_t R_Q[kWarpTileSeqLenQ][4];  // [1][4]
   uint32_t R_K[kWarpTileSeqLenK][2];  // [8][2]
diff --git a/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qk_swizzle_qk.cu b/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qk_swizzle_qk.cu
index abdedf24..8b1e09c5 100644
--- a/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qk_swizzle_qk.cu
+++ b/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qk_swizzle_qk.cu
@@ -324,16 +324,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d].
   uint32_t R_Q[kWarpTileSeqLenQ][4];  // [1][4]
   uint32_t R_K[kWarpTileSeqLenK][2];  // [8][2]
diff --git a/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qk_swizzle_qkv.cu b/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qk_swizzle_qkv.cu
index ace505e4..4cf8c460 100644
--- a/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qk_swizzle_qkv.cu
+++ b/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qk_swizzle_qkv.cu
@@ -328,16 +328,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d].
   uint32_t R_Q[kWarpTileSeqLenQ][4];  // [1][4]
   uint32_t R_K[kWarpTileSeqLenK][2];  // [8][2]
diff --git a/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_q.cu b/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_q.cu
index c127c7bf..b21db5e2 100644
--- a/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_q.cu
+++ b/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_q.cu
@@ -307,16 +307,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d].
   uint32_t R_Q[kWarpTileSeqLenQ][4]; // [1][4]
   uint32_t R_K[kWarpTileSeqLenK][2]; // [8][2]
diff --git a/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_q_F32F16F16F32.cu b/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_q_F32F16F16F32.cu
index 4b6ac6af..68a26e52 100644
--- a/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_q_F32F16F16F32.cu
+++ b/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_q_F32F16F16F32.cu
@@ -275,16 +275,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d].
   uint32_t R_Q[kWarpTileSeqLenQ][4]; // [1][4]
   uint32_t R_K[kWarpTileSeqLenK][2]; // [8][2]
diff --git a/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_qk.cu b/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_qk.cu
index 67f74012..5c1a67d0 100644
--- a/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_qk.cu
+++ b/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_qk.cu
@@ -307,16 +307,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d].
   uint32_t R_Q[kWarpTileSeqLenQ][4]; // [1][4]
   uint32_t R_K[kWarpTileSeqLenK][2]; // [8][2]
diff --git a/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_qk_F32F16F16F32.cu b/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_qk_F32F16F16F32.cu
index 7916d367..8c2ca6a0 100644
--- a/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_qk_F32F16F16F32.cu
+++ b/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_qk_F32F16F16F32.cu
@@ -275,16 +275,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d].
   uint32_t R_Q[kWarpTileSeqLenQ][4]; // [1][4]
   uint32_t R_K[kWarpTileSeqLenK][2]; // [8][2]
diff --git a/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_qkv.cu b/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_qkv.cu
index 276569ea..64c3074b 100644
--- a/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_qkv.cu
+++ b/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_qkv.cu
@@ -305,16 +305,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d].
   uint32_t R_Q[kWarpTileSeqLenQ][4]; // [1][4]
   uint32_t R_K[kWarpTileSeqLenK][2]; // [8][2]
diff --git a/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_qkv_F32F16F16F32.cu b/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_qkv_F32F16F16F32.cu
index 0141be86..50eb8af5 100644
--- a/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_qkv_F32F16F16F32.cu
+++ b/kernels/flash-attn/mma/swizzle/flash_attn_mma_tiling_qkv_swizzle_qkv_F32F16F16F32.cu
@@ -305,16 +305,16 @@ __global__ void __launch_bounds__(WARP_SIZE *kMmaTileSeqLenQ *kMmaTileSeqLenK)
   uint32_t smem_K_base_ptr = __cvta_generic_to_shared(K_tile_smem);
   uint32_t smem_V_base_ptr = __cvta_generic_to_shared(V_tile_smem);
 
-  // --------------------- Registers/SMEM for thread block
-  // ------------------------- block m_old, l_old, store in lane, use float to
+  // Registers/SMEM for thread block
+  // block m_old, l_old, store in lane, use float to
   // keep precision.
   float lane_block_row_max_old[kWarpTileSeqLenQ][2]; // [1][2]
   float lane_block_row_sum_old[kWarpTileSeqLenQ][2]; // [1][2]
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_max_old, -INFINITY);
   fill_2D_regs<float, kWarpTileSeqLenQ, 2>(lane_block_row_sum_old, 0.0f);
 
-  // ---------------------- Registers for S=Q@K^T/O=P@V
-  // ---------------------------- registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
+  // Registers for S=Q@K^T/O=P@V
+  // registers for QKV, S=Q[Br,d]@K[Bc,d]=[Br,Bc]
   // and O=P[Br,Bc]@V[Bc,d]=[Br,d].
   uint32_t R_Q[kWarpTileSeqLenQ][4]; // [1][4]
   uint32_t R_K[kWarpTileSeqLenK][2]; // [8][2]

From f819680960c7232f9d0fb0eb7d28dffba6436070 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Thu, 10 Jul 2025 20:07:58 +0800
Subject: [PATCH 39/57] Update README.md

---
 README.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/README.md b/README.md
index 75bbd459..67576b07 100644
--- a/README.md
+++ b/README.md
@@ -38,7 +38,7 @@
 ## 📖 News 🔥🔥
 <div id="news"></div>
 
-- [2025-06-16]: [🤗CacheDiT](https://github.com/vipshop/cache-dit) is release! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for Diffusion Transformers (**DBCache**, **DBPrune**, **FBCache**, etc)🔥. Feel free to take a try!
+- [2025-06-16]: [🤗CacheDiT](https://github.com/vipshop/cache-dit) is release! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for Diffusion Transformers (**DBCache**, **DBPrune**, **TaylorSeer**, **FBCache**, etc)🔥. Feel free to take a try!
 
 <div align='center'>
   <img src='https://github.com/user-attachments/assets/a5ec4320-d2f9-4254-888a-170b2d9e3784' height=170px>

From 0bcaeaa8c8a2898b12805fe9ebe482a9602b65b1 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Sun, 13 Jul 2025 17:38:34 +0800
Subject: [PATCH 40/57] Update README.md

---
 README.md | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/README.md b/README.md
index 67576b07..133f4caf 100644
--- a/README.md
+++ b/README.md
@@ -38,6 +38,8 @@
 ## 📖 News 🔥🔥
 <div id="news"></div>
 
+- [2025-07-13] An end2end speedup example for FLUX using cache-dit is released! **[🤗flux-faster](https://github.com/xlite-dev/flux-faster)**: A forked version of [huggingface/flux-fast](https://github.com/huggingface/flux-fast) that **makes flux-fast even faster** with **[cache-dit](https://github.com/vipshop/cache-dit)**, **3.3x** speedup on NVIDIA L20 while still maintaining **high precision**.
+  
 - [2025-06-16]: [🤗CacheDiT](https://github.com/vipshop/cache-dit) is release! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for Diffusion Transformers (**DBCache**, **DBPrune**, **TaylorSeer**, **FBCache**, etc)🔥. Feel free to take a try!
 
 <div align='center'>

From 89195986229f34be97259235b5fbbdccc0b4d2fa Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Mon, 14 Jul 2025 10:39:16 +0800
Subject: [PATCH 41/57] Update README.md

---
 README.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/README.md b/README.md
index 133f4caf..07ed1bd4 100644
--- a/README.md
+++ b/README.md
@@ -38,7 +38,7 @@
 ## 📖 News 🔥🔥
 <div id="news"></div>
 
-- [2025-07-13] An end2end speedup example for FLUX using cache-dit is released! **[🤗flux-faster](https://github.com/xlite-dev/flux-faster)**: A forked version of [huggingface/flux-fast](https://github.com/huggingface/flux-fast) that **makes flux-fast even faster** with **[cache-dit](https://github.com/vipshop/cache-dit)**, **3.3x** speedup on NVIDIA L20 while still maintaining **high precision**.
+- [2025-07-13] **[🤗flux-faster](https://github.com/xlite-dev/flux-faster)** is released! A forked version of [huggingface/flux-fast](https://github.com/huggingface/flux-fast) that **makes flux-fast even faster** with **[cache-dit](https://github.com/vipshop/cache-dit)**, **3.3x** speedup on NVIDIA L20 while still maintaining **high precision**.
   
 - [2025-06-16]: [🤗CacheDiT](https://github.com/vipshop/cache-dit) is release! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for Diffusion Transformers (**DBCache**, **DBPrune**, **TaylorSeer**, **FBCache**, etc)🔥. Feel free to take a try!
 

From cb6a0497e829741fc5e3ad5e43a5beed0d85edcd Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Wed, 16 Jul 2025 21:12:14 +0800
Subject: [PATCH 42/57] Update README.md

---
 README.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/README.md b/README.md
index 07ed1bd4..d1b7e46a 100644
--- a/README.md
+++ b/README.md
@@ -40,7 +40,7 @@
 
 - [2025-07-13] **[🤗flux-faster](https://github.com/xlite-dev/flux-faster)** is released! A forked version of [huggingface/flux-fast](https://github.com/huggingface/flux-fast) that **makes flux-fast even faster** with **[cache-dit](https://github.com/vipshop/cache-dit)**, **3.3x** speedup on NVIDIA L20 while still maintaining **high precision**.
   
-- [2025-06-16]: [🤗CacheDiT](https://github.com/vipshop/cache-dit) is release! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for Diffusion Transformers (**DBCache**, **DBPrune**, **TaylorSeer**, **FBCache**, etc)🔥. Feel free to take a try!
+- [2025-06-16]: [🤗CacheDiT](https://github.com/vipshop/cache-dit) is released! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for Diffusion Transformers (**DBCache**, **DBPrune**, **TaylorSeer**, **FBCache**, etc)🔥. Feel free to take a try!
 
 <div align='center'>
   <img src='https://github.com/user-attachments/assets/a5ec4320-d2f9-4254-888a-170b2d9e3784' height=170px>

From 2f75697fd01f69f6381fde70ab31ac6bcbcdb50c Mon Sep 17 00:00:00 2001
From: kitecats <90194592+kitecats@users.noreply.github.com>
Date: Mon, 21 Jul 2025 16:35:06 +0800
Subject: [PATCH 43/57] feat: Add a naive ws-hgemm for sm8x (#366)

* Use 128-bit data loading

* add a naive ws-hgemm in sm8x

* rename wshgemm to ws-hgemm
---
 README.md                               |   2 +-
 kernels/layer-norm/layer_norm.py        |   2 +-
 kernels/ws-hgemm/README.md              |  23 ++
 kernels/ws-hgemm/naive_ws_hgemm_sm8x.cu | 486 ++++++++++++++++++++++++
 kernels/ws-hgemm/ws_hgemm.py            |  81 ++++
 5 files changed, 592 insertions(+), 2 deletions(-)
 create mode 100755 kernels/ws-hgemm/README.md
 create mode 100644 kernels/ws-hgemm/naive_ws_hgemm_sm8x.cu
 create mode 100644 kernels/ws-hgemm/ws_hgemm.py

diff --git a/README.md b/README.md
index d1b7e46a..462b561c 100644
--- a/README.md
+++ b/README.md
@@ -39,7 +39,7 @@
 <div id="news"></div>
 
 - [2025-07-13] **[🤗flux-faster](https://github.com/xlite-dev/flux-faster)** is released! A forked version of [huggingface/flux-fast](https://github.com/huggingface/flux-fast) that **makes flux-fast even faster** with **[cache-dit](https://github.com/vipshop/cache-dit)**, **3.3x** speedup on NVIDIA L20 while still maintaining **high precision**.
-  
+
 - [2025-06-16]: [🤗CacheDiT](https://github.com/vipshop/cache-dit) is released! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for Diffusion Transformers (**DBCache**, **DBPrune**, **TaylorSeer**, **FBCache**, etc)🔥. Feel free to take a try!
 
 <div align='center'>
diff --git a/kernels/layer-norm/layer_norm.py b/kernels/layer-norm/layer_norm.py
index e9c68400..8cb94e77 100644
--- a/kernels/layer-norm/layer_norm.py
+++ b/kernels/layer-norm/layer_norm.py
@@ -98,7 +98,7 @@ def run_benchmark(
 
 print(" " * 40 + f"f16 overflow without f32")
 print("-" * 85)
-x_f16 = x.half() * 100 # this will cause overflow for kernels without `f32`
+x_f16 = x.half() * 100  # this will cause overflow for kernels without `f32`
 run_benchmark(lib.layer_norm_f16_f16, x_f16, "f16f16", out_f16)
 run_benchmark(lib.layer_norm_f16_f32, x_f16, "f16f32", out_f16)
 run_benchmark(lib.layer_norm_f16x2_f16, x_f16, "f16x2f16", out_f16)
diff --git a/kernels/ws-hgemm/README.md b/kernels/ws-hgemm/README.md
new file mode 100755
index 00000000..fde79129
--- /dev/null
+++ b/kernels/ws-hgemm/README.md
@@ -0,0 +1,23 @@
+# Warp Specialization HGemm
+
+## 0x00 说明
+
+包含以下内容：
+
+- [X] ws_hgemm_naive_cute_kernel
+
+
+## 测试
+
+```bash
+python3 ws_hgemm.py
+```
+
+输出:
+
+```bash
+--------------------------------------------------------------------------------
+out_ws_hgemm_naive_cute: [4096.0, 4096.0, 4096.0], time:3.71974587ms
+out_f16_th: [4096.0, 4096.0, 4096.0], time:5.05561471ms
+--------------------------------------------------------------------------------
+```
diff --git a/kernels/ws-hgemm/naive_ws_hgemm_sm8x.cu b/kernels/ws-hgemm/naive_ws_hgemm_sm8x.cu
new file mode 100644
index 00000000..5a28febd
--- /dev/null
+++ b/kernels/ws-hgemm/naive_ws_hgemm_sm8x.cu
@@ -0,0 +1,486 @@
+﻿#pragma once
+#include "cooperative_groups.h"
+#include "cuda/pipeline"
+#include "cute/tensor.hpp"
+#include <cooperative_groups/memcpy_async.h>
+#include <torch/extension.h>
+#include <torch/types.h>
+
+#define DEVICE __device__ __forceinline__
+
+using namespace cute;
+
+template <class CTATile, int ProducerThread, int Stage> struct WSHGEMMTraits {
+  using MatrixTypeAB = half;
+  using AccType = half;
+
+  constexpr static int kCTAM = get<0>(CTATile{});
+  constexpr static int kCTAN = get<1>(CTATile{});
+  constexpr static int kCTAK = get<2>(CTATile{});
+  constexpr static int kStage = Stage;
+
+  // MMA Trait
+  using mma_op = SM80_16x8x16_F16F16F16F16_TN;
+  using mma_traits = MMA_Traits<mma_op>;
+  using mma_atom = MMA_Atom<mma_traits>;
+  using mma_atom_shape = mma_traits::Shape_MNK;
+
+  constexpr static int kMmaThrLayoutM = 2;
+  constexpr static int kMmaThrLayoutN = 2;
+  constexpr static int kMmaThrLayoutK = 1;
+
+  constexpr static int kMmaPermuteM = kMmaThrLayoutM * get<0>(mma_atom_shape{});
+  constexpr static int kMmaPermuteN =
+      2 * kMmaThrLayoutN * get<1>(mma_atom_shape{});
+  constexpr static int kMmaPermuteK = kMmaThrLayoutK * get<2>(mma_atom_shape{});
+
+  using MmaThrLayout = decltype(make_layout(make_shape(
+      Int<kMmaThrLayoutM>{}, Int<kMmaThrLayoutN>{}, Int<kMmaThrLayoutK>{})));
+
+  using MmaPermutation = decltype(make_tile(
+      Int<kMmaPermuteM>{}, Int<kMmaPermuteN>{}, Int<kMmaPermuteK>{}));
+
+  // The expanded TiledMMA can process matrices of size 32x32x16 in a single
+  // operation.
+  using TiledMMA =
+      decltype(make_tiled_mma(mma_atom{}, MmaThrLayout{}, MmaPermutation{}));
+
+  constexpr static int kConsumerThread = size(TiledMMA{});
+  // To avoid warp divergence
+  static_assert(ProducerThread % 32 == 0,
+                "The number of ProducerThreads must be a multiple of 32");
+  constexpr static int kProducerThread = ProducerThread;
+  constexpr static int kAllThread = kProducerThread + kConsumerThread;
+
+  // Smem
+  constexpr static int kSwizzleB = 3;
+  constexpr static int kSwizzleM = 3;
+  constexpr static int kSwizzleS = 3;
+
+  using SmemLayoutAtom =
+      decltype(composition(Swizzle<kSwizzleB, kSwizzleM, kSwizzleS>{},
+                           make_layout(make_shape(Int<8>{}, Int<kCTAK>{}),
+                                       make_stride(Int<kCTAK>{}, Int<1>{}))));
+
+  using SmemLayoutA = decltype(tile_to_shape(
+      SmemLayoutAtom{}, make_shape(Int<kCTAM>{}, Int<kCTAK>{}, Int<kStage>{})));
+
+  using SmemLayoutB = decltype(tile_to_shape(
+      SmemLayoutAtom{}, make_shape(Int<kCTAN>{}, Int<kCTAK>{}, Int<kStage>{})));
+
+  constexpr static int kSmemSizeA = cosize(SmemLayoutA{});
+  constexpr static int kSmemSizeB = cosize(SmemLayoutB{});
+  constexpr static int kSmemAllocateAB =
+      (kSmemSizeA + kSmemSizeB) * sizeof(MatrixTypeAB);
+
+  constexpr static int kSmemStageAcc = 2;
+  using SmemLayoutAcc = decltype(tile_to_shape(
+      SmemLayoutAtom{}, make_shape(Int<kMmaPermuteM>{}, Int<kMmaPermuteN>{},
+                                   Int<kSmemStageAcc>{})));
+
+  constexpr static int kSmemSizeAcc = cosize(SmemLayoutAcc{});
+  constexpr static int kSmemAllocateAcc = kSmemSizeAcc * sizeof(AccType);
+  constexpr static int kAllSmemAllocate =
+      cute::max(kSmemAllocateAB, kSmemAllocateAcc);
+
+  // Producer g2s copy
+  using g2s_copy_op = SM80_CP_ASYNC_CACHEGLOBAL<cute::uint128_t>;
+  using g2s_copy_traits = Copy_Traits<g2s_copy_op>;
+  using g2s_copy_atom = Copy_Atom<g2s_copy_traits, MatrixTypeAB>;
+
+  constexpr static int g2s_thread_vec_size =
+      sizeof(cute::uint128_t) / sizeof(MatrixTypeAB);
+  constexpr static int g2s_thread_k = kCTAK / g2s_thread_vec_size;
+  constexpr static int g2s_thread_m = kProducerThread / g2s_thread_k;
+
+  using G2SCopyA = decltype(make_tiled_copy(
+      g2s_copy_atom{},
+      make_layout(make_shape(Int<g2s_thread_m>{}, Int<g2s_thread_k>{}),
+                  make_stride(Int<g2s_thread_k>{}, Int<1>{})),
+      make_layout(make_shape(Int<1>{}, Int<g2s_thread_vec_size>{}))));
+
+  using G2SCopyB = G2SCopyA;
+
+  // Consumer s2r copy
+  using s2r_copy_op = SM75_U32x4_LDSM_N;
+  using s2r_copy_traits = Copy_Traits<s2r_copy_op>;
+  using s2r_copy_atom = Copy_Atom<s2r_copy_traits, MatrixTypeAB>;
+
+  using S2RCopyA = s2r_copy_atom;
+  using S2RCopyB = s2r_copy_atom;
+
+  // Consumer r2s copy
+  using R2SCopyC = Copy_Atom<UniversalCopy<int>, AccType>;
+
+  // Consumer s2g copy
+  using S2GCopyAtomC = Copy_Atom<UniversalCopy<cute::uint128_t>, AccType>;
+  constexpr static int s2g_thread_vec_size =
+      sizeof(cute::uint128_t) / sizeof(AccType);
+  constexpr static int s2g_thread_n = kMmaPermuteN / s2g_thread_vec_size;
+  constexpr static int s2g_thread_m = kConsumerThread / s2g_thread_n;
+
+  using S2GCopyC = decltype(make_tiled_copy(
+      S2GCopyAtomC{},
+      make_layout(make_shape(Int<s2g_thread_m>{}, Int<s2g_thread_n>{}),
+                  make_stride(Int<s2g_thread_n>{}, Int<1>{})),
+      make_layout(make_shape(Int<1>{}, Int<s2g_thread_vec_size>{}))));
+
+  struct Arguments {
+    void *a_ptr;
+    void *b_ptr;
+    void *c_ptr;
+
+    using Gemm_Shape = Shape<int, int, int>;
+
+    Gemm_Shape problem_shape;
+    Gemm_Shape tile_shape;
+
+    Arguments() = delete;
+    Arguments(Gemm_Shape problem_shape_, void *a_ptr_, void *b_ptr_,
+              void *c_ptr_)
+        : problem_shape(problem_shape_), a_ptr(a_ptr_), b_ptr(b_ptr_),
+          c_ptr(c_ptr_) {
+      int m = size<0>(problem_shape);
+      int n = size<1>(problem_shape);
+      int k = size<2>(problem_shape);
+
+      int m_tiles = ceil_div(m, kCTAM);
+      int n_tiles = ceil_div(n, kCTAN);
+      int k_tiles = ceil_div(k, kCTAK);
+      tile_shape = make_shape(m_tiles, n_tiles, k_tiles);
+    }
+
+    dim3 get_grid() {
+      return dim3(ceil_div(get<0>(problem_shape), kCTAM),
+                  ceil_div(get<1>(problem_shape), kCTAN));
+    }
+  };
+
+  template <typename Pipeline, typename AEngine, typename ALayout,
+            typename BEngine, typename BLayout>
+  DEVICE static auto producer(void *smem_ptr, Pipeline &pipeline,
+                              Tensor<AEngine, ALayout> const &gA,
+                              Tensor<BEngine, BLayout> const &gB) {
+    using T = typename WSHGEMMTraits::MatrixTypeAB;
+    constexpr int SmemSizeA = WSHGEMMTraits::kSmemSizeA;
+    auto tidx = threadIdx.x;
+
+    T *SmemPtrA = reinterpret_cast<T *>(smem_ptr);
+    T *SmemPtrB = SmemPtrA + SmemSizeA;
+
+    auto sA = make_tensor(make_smem_ptr<T>(SmemPtrA),
+                          SmemLayoutA{}); // (CTAM, CTAK, Stage)
+    auto sB = make_tensor(make_smem_ptr<T>(SmemPtrB),
+                          SmemLayoutB{}); // (CTAN, CTAK, Stage)
+
+    G2SCopyA g2s_copy_A;
+    auto thr_g2s_copy_A = g2s_copy_A.get_slice(tidx);
+    auto g2s_tAgA = thr_g2s_copy_A.partition_S(gA);
+    auto g2s_tAsA = thr_g2s_copy_A.partition_D(sA);
+
+    G2SCopyB g2s_copy_B;
+    auto thr_g2s_copy_B = g2s_copy_B.get_slice(tidx);
+    auto g2s_tBgB = thr_g2s_copy_B.partition_S(gB);
+    auto g2s_tBsB = thr_g2s_copy_B.partition_D(sB);
+
+    const auto kNumIterationK = size<2>(gA);
+    int g2s_g_read_idx = 0;
+    int g2s_s_write_idx = 0;
+
+    for (int iter_k = 0, stage_idx = 0; iter_k < kNumIterationK; iter_k++) {
+      for (; stage_idx < kNumIterationK && stage_idx < (iter_k + kStage);
+           stage_idx++) {
+
+        pipeline.producer_acquire();
+
+        copy(g2s_copy_A, g2s_tAgA(_, _, _, g2s_g_read_idx),
+             g2s_tAsA(_, _, _, g2s_s_write_idx));
+
+        copy(g2s_copy_B, g2s_tBgB(_, _, _, g2s_g_read_idx),
+             g2s_tBsB(_, _, _, g2s_s_write_idx));
+
+        pipeline.producer_commit();
+
+        g2s_g_read_idx++;
+        g2s_s_write_idx = (g2s_s_write_idx + 1) % kStage;
+      }
+    }
+  }
+
+  template <typename Pipeline, typename CEngine, typename CLayout>
+  DEVICE static auto main_loop(Arguments const &args, void *smem_ptr,
+                               Pipeline &pipeline,
+                               Tensor<CEngine, CLayout> const &gC) {
+    using T = typename WSHGEMMTraits::MatrixTypeAB;
+    constexpr int SmemSizeA = WSHGEMMTraits::kSmemSizeA;
+    auto tidx = threadIdx.x - kProducerThread;
+
+    T *SmemPtrA = reinterpret_cast<T *>(smem_ptr);
+    T *SmemPtrB = SmemPtrA + SmemSizeA;
+
+    auto sA = make_tensor(make_smem_ptr<T>(SmemPtrA),
+                          SmemLayoutA{}); // (CTAM, CTAK, Stage)
+    auto sB = make_tensor(make_smem_ptr<T>(SmemPtrB),
+                          SmemLayoutB{}); // (CTAN, CTAK, Stage)
+
+    TiledMMA tiled_mma;
+    auto thr_mma = tiled_mma.get_slice(tidx);
+
+    auto tArA = thr_mma.partition_fragment_A(sA(_, _, 0));
+    auto tBrB = thr_mma.partition_fragment_B(sB(_, _, 0));
+    auto tCrC = thr_mma.partition_fragment_C(gC);
+    clear(tCrC);
+
+    // s2r
+    auto s2r_copy_A = make_tiled_copy_A(S2RCopyA{}, tiled_mma);
+    auto thr_s2r_copy_A = s2r_copy_A.get_slice(tidx);
+    auto s2r_tAsA = thr_s2r_copy_A.partition_S(sA);
+    auto s2r_tArA_view = thr_s2r_copy_A.retile_D(tArA);
+
+    auto s2r_copy_B = make_tiled_copy_B(S2RCopyB{}, tiled_mma);
+    auto thr_s2r_copy_B = s2r_copy_B.get_slice(tidx);
+    auto s2r_tBsB = thr_s2r_copy_B.partition_S(sB);
+    auto s2r_tBrB_view = thr_s2r_copy_B.retile_D(tBrB);
+
+    const int kNumIterationK = get<2>(args.tile_shape);
+    const int kNumInnerStage = size<2>(tArA);
+
+    int s2r_s_read_idx = 0;
+    int next_s2r_s_read_idx = 0;
+
+    if (kNumInnerStage > 1) {
+      pipeline.consumer_wait();
+
+      copy(s2r_copy_A, s2r_tAsA(_, _, 0, s2r_s_read_idx),
+           s2r_tArA_view(_, _, 0));
+      copy(s2r_copy_B, s2r_tBsB(_, _, 0, s2r_s_read_idx),
+           s2r_tBrB_view(_, _, 0));
+    }
+
+    for (int iter_k = 0; iter_k < kNumIterationK; iter_k++) {
+#pragma unroll
+      for (int inner_stage = 0; inner_stage < kNumInnerStage; inner_stage++) {
+        int next_inner_stage = (inner_stage + 1) % kNumInnerStage;
+
+        if (inner_stage == kNumInnerStage - 1 && iter_k < kNumIterationK - 1) {
+          pipeline.consumer_wait();
+          s2r_s_read_idx = next_s2r_s_read_idx;
+        }
+
+        copy(s2r_copy_A, s2r_tAsA(_, _, next_inner_stage, s2r_s_read_idx),
+             s2r_tArA_view(_, _, next_inner_stage));
+
+        copy(s2r_copy_B, s2r_tBsB(_, _, next_inner_stage, s2r_s_read_idx),
+             s2r_tBrB_view(_, _, next_inner_stage));
+
+        if (inner_stage == 0) {
+          pipeline.consumer_release();
+          next_s2r_s_read_idx = (s2r_s_read_idx + 1) % kStage;
+        }
+
+        gemm(tiled_mma, tArA(_, _, inner_stage), tBrB(_, _, inner_stage), tCrC);
+      }
+    }
+
+    return tCrC;
+  }
+
+  template <typename AccEngine, typename AccLayout, typename CEngine,
+            typename CLayout>
+  DEVICE static void epilog(Arguments const &args, void *smem_ptr,
+                            Tensor<AccEngine, AccLayout> const &acc,
+                            Tensor<CEngine, CLayout> &gC) {
+    __syncthreads(); // wait all consumer thread finish main_loop
+
+    int tidx = threadIdx.x - kProducerThread;
+
+    auto sC = make_tensor(
+        make_smem_ptr<AccType>(smem_ptr),
+        SmemLayoutAcc{}); // (kMmaPermuteM, kMmaPermuteN, kSmemStageAcc)
+
+    // r2s
+    auto r2s_copy_C = make_tiled_copy_C(R2SCopyC{}, TiledMMA{});
+    auto thr_r2s_copy_C = r2s_copy_C.get_slice(tidx);
+    auto r2s_tCrC = thr_r2s_copy_C.retile_S(acc);
+    auto r2s_tCsC = thr_r2s_copy_C.partition_D(sC);
+
+    auto r2s_tCrC_view = group_modes<1, 3>(r2s_tCrC);
+
+    // s2g
+    S2GCopyC s2g_copy_C;
+    auto thr_s2g_copy_C = s2g_copy_C.get_slice(tidx);
+    auto s2g_tCsC = thr_s2g_copy_C.partition_S(sC);
+    auto s2g_tCgC = thr_s2g_copy_C.partition_D(gC);
+
+    auto s2g_tCgC_view = group_modes<1, 3>(s2g_tCgC);
+
+    const int kEpilogIterations = size<1>(r2s_tCrC_view);
+    const int kEpilogStages = size<3>(r2s_tCsC);
+
+#pragma unroll
+    for (int epilog_iter = 0; epilog_iter < kEpilogIterations;
+         epilog_iter += kEpilogStages) {
+#pragma unroll
+      for (int epilog_stage_idx = 0; epilog_stage_idx < kEpilogStages;
+           epilog_stage_idx++) {
+        // r2s
+        copy(r2s_copy_C, r2s_tCrC_view(_, epilog_iter + epilog_stage_idx),
+             r2s_tCsC(_, 0, 0, epilog_stage_idx));
+      }
+      __syncthreads(); // wait all consumer thread finish r2s
+
+#pragma unroll
+      for (int epilog_stage_idx = 0; epilog_stage_idx < kEpilogStages;
+           epilog_stage_idx++) {
+        // s2g
+        copy(s2g_copy_C, s2g_tCsC(_, 0, 0, epilog_stage_idx),
+             s2g_tCgC_view(_, epilog_iter + epilog_stage_idx));
+      }
+      __syncthreads(); // wait all consumer thread finish s2g
+    }
+  }
+
+  template <typename Pipeline, typename CEngine, typename CLayout>
+  DEVICE static void consumer(Arguments const &args, void *smem_ptr,
+                              Pipeline &pipeline,
+                              Tensor<CEngine, CLayout> &gC) {
+    auto tCrC = main_loop(args, smem_ptr, pipeline, gC);
+    epilog(args, smem_ptr, tCrC, gC);
+  }
+};
+
+#pragma nv_diag_suppress static_var_with_dynamic_init
+template <typename WSHGEMMTraits>
+__global__ void
+ws_hgemm_naive_cute_kernel(typename WSHGEMMTraits::Arguments args) {
+  using MatrixTypeAB = typename WSHGEMMTraits::MatrixTypeAB;
+  using AccType = typename WSHGEMMTraits::AccType;
+
+  constexpr int kCTAM = WSHGEMMTraits::kCTAM;
+  constexpr int kCTAN = WSHGEMMTraits::kCTAN;
+  constexpr int kCTAK = WSHGEMMTraits::kCTAK;
+  constexpr int kStage = WSHGEMMTraits::kStage;
+
+  auto block = cooperative_groups::this_thread_block();
+  auto tidx = threadIdx.x;
+
+  auto tile_id_m = blockIdx.x;
+  auto tile_id_n = blockIdx.y;
+
+  // set thread role
+  const auto thread_role = tidx < WSHGEMMTraits::kProducerThread
+                               ? cuda::pipeline_role::producer
+                               : cuda::pipeline_role::consumer;
+  // create pipeline
+  __shared__ cuda::pipeline_shared_state<cuda::thread_scope::thread_scope_block,
+                                         kStage>
+      shared_state;
+  auto pipeline = cuda::make_pipeline(block, &shared_state, thread_role);
+
+  extern __shared__ MatrixTypeAB smem_ptr[];
+
+  auto A =
+      make_tensor(make_gmem_ptr<MatrixTypeAB>(args.a_ptr),
+                  select<0, 2>(args.problem_shape), GenRowMajor{}); // (M, K)
+
+  auto B =
+      make_tensor(make_gmem_ptr<MatrixTypeAB>(args.b_ptr),
+                  select<1, 2>(args.problem_shape), GenRowMajor{}); // (N, K)
+
+  auto C =
+      make_tensor(make_gmem_ptr<AccType>(args.c_ptr),
+                  select<0, 1>(args.problem_shape), GenRowMajor{}); // (M, N)
+
+  auto gA = local_tile(A, make_tile(Int<kCTAM>{}, Int<kCTAK>{}),
+                       make_coord(tile_id_m, _)); // (kCTAM, kCTAK, K/kCTAK)
+
+  auto gB = local_tile(B, make_tile(Int<kCTAN>{}, Int<kCTAK>{}),
+                       make_coord(tile_id_n, _)); // (kCTAN, kCTAK, K/kCTAK)
+
+  auto gC = local_tile(C, make_tile(Int<kCTAM>{}, Int<kCTAN>{}),
+                       make_coord(tile_id_m, tile_id_n)); // (kCTAM, kCTAN)
+
+  // Different thread_roles execute different branches.
+  if (thread_role == cuda::pipeline_role::producer) {
+    WSHGEMMTraits::producer(smem_ptr, pipeline, gA, gB);
+  } else {
+    WSHGEMMTraits::consumer(args, smem_ptr, pipeline, gC);
+  }
+}
+
+#define STRINGFY(str) #str
+#define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
+  m.def(STRINGFY(func), &func, STRINGFY(func));
+
+#define CHECK_TORCH_TENSOR_DTYPE(T, th_type)                                   \
+  if (((T).options().dtype() != (th_type))) {                                  \
+    std::cout << "Tensor Info:" << (T).options() << std::endl;                 \
+    throw std::runtime_error("values must be " #th_type);                      \
+  }
+
+#define CHECK_TORCH_TENSOR_SHAPE(T, S0, S1)                                    \
+  if (((T).size(0) != (S0)) || ((T).size(1) != (S1))) {                        \
+    throw std::runtime_error("Tensor size mismatch!");                         \
+  }
+
+inline int get_max_smem_size() {
+  int max_shared_mem;
+  cudaDeviceGetAttribute(&max_shared_mem,
+                         cudaDevAttrMaxSharedMemoryPerBlockOptin, 0);
+  return max_shared_mem;
+}
+
+template <typename Kernel> void config_smem(Kernel kernel, int smem_size) {
+  if (smem_size >= 32 * 1024) {
+    if (cudaFuncSetAttribute(kernel,
+                             cudaFuncAttributeMaxDynamicSharedMemorySize,
+                             smem_size) != cudaSuccess) {
+      int max_shared_mem = get_max_smem_size();
+      cudaError_t err = cudaGetLastError();
+      std::cerr << "Set kernel attribute failed: " << cudaGetErrorString(err)
+                << std::endl;
+      std::cerr
+          << "Kernel required " << smem_size
+          << " shared memory but the max shared memory per block optin is: "
+          << max_shared_mem << std::endl;
+    }
+  }
+}
+
+// WarpSpecialization HGEMM
+void ws_hgemm_naive_cute(torch::Tensor a, torch::Tensor b, torch::Tensor c) {
+
+  CHECK_TORCH_TENSOR_DTYPE(a, torch::kHalf)
+  CHECK_TORCH_TENSOR_DTYPE(b, torch::kHalf)
+  CHECK_TORCH_TENSOR_DTYPE(c, torch::kHalf)
+  const int M = a.size(0);
+  const int K = a.size(1);
+  const int N = b.size(1);
+  CHECK_TORCH_TENSOR_SHAPE(a, M, K)
+  CHECK_TORCH_TENSOR_SHAPE(b, K, N)
+  CHECK_TORCH_TENSOR_SHAPE(c, M, N)
+
+  using GEMM_Traits =
+      WSHGEMMTraits<decltype(make_shape(_128{}, _256{}, _32{})), 32, 3>;
+
+  // set smem size
+  constexpr int smem_size = GEMM_Traits::kAllSmemAllocate;
+  config_smem(ws_hgemm_naive_cute_kernel<GEMM_Traits>, smem_size);
+
+  // set problem size
+  using Arguments = typename GEMM_Traits::Arguments;
+  Arguments args(make_shape(M, N, K), reinterpret_cast<half *>(a.data_ptr()),
+                 reinterpret_cast<half *>(b.data_ptr()),
+                 reinterpret_cast<half *>(c.data_ptr()));
+
+  constexpr int block_size = GEMM_Traits::kAllThread;
+
+  dim3 block(block_size);
+  dim3 grid(args.get_grid());
+
+  ws_hgemm_naive_cute_kernel<GEMM_Traits><<<grid, block, smem_size>>>(args);
+}
+
+PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
+  TORCH_BINDING_COMMON_EXTENSION(ws_hgemm_naive_cute)
+}
diff --git a/kernels/ws-hgemm/ws_hgemm.py b/kernels/ws-hgemm/ws_hgemm.py
new file mode 100644
index 00000000..678cf1ab
--- /dev/null
+++ b/kernels/ws-hgemm/ws_hgemm.py
@@ -0,0 +1,81 @@
+﻿import os
+import time
+from functools import partial
+from typing import Optional
+
+import torch
+from torch.utils.cpp_extension import load
+
+torch.set_grad_enabled(False)
+
+CUTLASS_REPO_PATH = os.environ.get(
+    "CUTLASS_REPO_PATH", os.path.expanduser("../../third-party/cutlass")
+)
+
+# Load the CUDA kernel as a python module
+lib = load(
+    name="naive_ws_hgemm_sm8x_lib",
+    sources=["naive_ws_hgemm_sm8x.cu"],
+    extra_cuda_cflags=[
+        "-O3",
+        "-U__CUDA_NO_HALF_OPERATORS__",
+        "-U__CUDA_NO_HALF_CONVERSIONS__",
+        "-U__CUDA_NO_HALF2_OPERATORS__",
+        "-U__CUDA_NO_BFLOAT16_CONVERSIONS__",
+        "--expt-relaxed-constexpr",
+        "--expt-extended-lambda",
+        "--use_fast_math",
+    ],
+    extra_cflags=["-std=c++17"],
+    extra_include_paths=[os.path.join(CUTLASS_REPO_PATH, "include")],
+)
+
+
+def run_benchmark(
+    perf_func: callable,
+    a: torch.Tensor,
+    b: torch.Tensor,
+    tag: str,
+    out: Optional[torch.Tensor] = None,
+    warmup: int = 10,
+    iters: int = 200,
+    show_all: bool = False,
+):
+    if out is not None:
+        out.fill_(0)
+    if out is not None:
+        for i in range(warmup):
+            perf_func(a, b, out)
+    else:
+        for i in range(warmup):
+            _ = perf_func(a, b)
+
+    torch.cuda.synchronize()
+    start = time.time()
+    # iters
+    if out is not None:
+        for i in range(iters):
+            perf_func(a, b, out)
+    else:
+        for i in range(iters):
+            out = perf_func(a, b)
+    torch.cuda.synchronize()
+    end = time.time()
+    total_time = (end - start) * 1000  # ms
+    mean_time = total_time / iters
+    out_info = f"out_{tag}"
+    out_val = out.flatten().detach().cpu().numpy().tolist()[:3]
+    out_val = [round(v, 8) for v in out_val]
+    print(f"{out_info:>13}: {out_val}, time:{mean_time:.8f}ms")
+    if show_all:
+        print(out)
+    return out.clone(), mean_time
+
+
+print("-" * 80)
+M, N, K = 4096, 4096, 4096
+a = torch.ones((M, K)).cuda().half().contiguous()
+b = torch.ones((K, N)).cuda().half().contiguous()
+c = torch.randn((M, N)).cuda().half().contiguous()
+run_benchmark(lib.ws_hgemm_naive_cute, a, b, "ws_hgemm_naive_cute", c)
+run_benchmark(partial(torch.matmul, out=c), a, b, "f16_th")

From 7676e177fd532adcff996f511d063f7d4e44ad69 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Wed, 30 Jul 2025 11:14:58 +0800
Subject: [PATCH 44/57] Update README.md

---
 README.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/README.md b/README.md
index 462b561c..70799983 100644
--- a/README.md
+++ b/README.md
@@ -40,7 +40,7 @@
 
 - [2025-07-13] **[🤗flux-faster](https://github.com/xlite-dev/flux-faster)** is released! A forked version of [huggingface/flux-fast](https://github.com/huggingface/flux-fast) that **makes flux-fast even faster** with **[cache-dit](https://github.com/vipshop/cache-dit)**, **3.3x** speedup on NVIDIA L20 while still maintaining **high precision**.
 
-- [2025-06-16]: [🤗CacheDiT](https://github.com/vipshop/cache-dit) is released! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for Diffusion Transformers (**DBCache**, **DBPrune**, **TaylorSeer**, **FBCache**, etc)🔥. Feel free to take a try!
+- [2025-06-16]: **[🤗cache-dit](https://github.com/vipshop/cache-dit)** is released! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for Diffusion Transformers (**DBCache**, **DBPrune**, **TaylorSeer**, **FBCache**, **Cache CFG**, etc)🔥. Feel free to take a try!
 
 <div align='center'>
   <img src='https://github.com/user-attachments/assets/a5ec4320-d2f9-4254-888a-170b2d9e3784' height=170px>

From 6a22ee791d2dc10fb2af3f2232c1c6d3121c6a86 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Wed, 30 Jul 2025 14:25:01 +0800
Subject: [PATCH 45/57] Update README.md

---
 README.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/README.md b/README.md
index 70799983..21241cfe 100644
--- a/README.md
+++ b/README.md
@@ -40,7 +40,7 @@
 
 - [2025-07-13] **[🤗flux-faster](https://github.com/xlite-dev/flux-faster)** is released! A forked version of [huggingface/flux-fast](https://github.com/huggingface/flux-fast) that **makes flux-fast even faster** with **[cache-dit](https://github.com/vipshop/cache-dit)**, **3.3x** speedup on NVIDIA L20 while still maintaining **high precision**.
 
-- [2025-06-16]: **[🤗cache-dit](https://github.com/vipshop/cache-dit)** is released! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for Diffusion Transformers (**DBCache**, **DBPrune**, **TaylorSeer**, **FBCache**, **Cache CFG**, etc)🔥. Feel free to take a try!
+- [2025-06-16]: **[🤗cache-dit](https://github.com/vipshop/cache-dit)** is released! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for DiTs (**DBCache**, **DBPrune**, **TaylorSeer**, **FBCache**, **Cache CFG**, etc)🔥. Feel free to take a try!
 
 <div align='center'>
   <img src='https://github.com/user-attachments/assets/a5ec4320-d2f9-4254-888a-170b2d9e3784' height=170px>

From 56e5e37efa763c5a33cba7bca53d16797f002189 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Fri, 1 Aug 2025 14:37:39 +0800
Subject: [PATCH 46/57] Add cute ws-hgemm index (#368)

Co-authored-by: yanjun.qiu <yanjun.qiu@vipshop.com>
---
 README.md | 1 +
 1 file changed, 1 insertion(+)

diff --git a/README.md b/README.md
index 21241cfe..81cf9243 100644
--- a/README.md
+++ b/README.md
@@ -477,6 +477,7 @@ The kernels listed here will guide you through a step-by-step progression, rangi
 | ✔️ [hgemv_f16x8_cute_kernel](./kernels/hgemv/hgemv_cute.cu)|f16|f16|[link](./kernels/hgemv/)|⭐️⭐️⭐️|
 | ✔️ [hgemv_tensor_core_cute_kernel](./kernels/hgemv/hgemv_cute.cu)|f16|f16|[link](./kernels/hgemv/)|⭐️⭐️⭐️|
 | ✔️ [hgemm_mma_stages_swizzle{smem}...cute*](./kernels/hgemm/cutlass/hgemm_mma_stage_tn_cute.cu)|f16|f16|[link](./kernels/hgemm/)|⭐️⭐️⭐️|
+| ✔️ [ws_hgemm_naive_cute_kernel](./kernels/ws-hgemm/naive_ws_hgemm_sm8x.cu)|f16|f16|[link](./kernels/ws-hgemm/)|⭐️⭐️⭐️|
 
 ## 📖 100+ 高性能计算与分布式-技术博客
 

From 1f9fb5d5de649ef03281609e1b3ce1c45f82359e Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Wed, 6 Aug 2025 17:45:30 +0800
Subject: [PATCH 47/57] misc: automated submodule update (#369)

---
 ffpa-attn           | 2 +-
 third-party/cutlass | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/ffpa-attn b/ffpa-attn
index 3b90cc1d..0b449617 160000
--- a/ffpa-attn
+++ b/ffpa-attn
@@ -1 +1 @@
-Subproject commit 3b90cc1d47eb921503dd57848e13a52f2a8cf2be
+Subproject commit 0b449617e8e45d3359131301d282de6bd4da8de0
diff --git a/third-party/cutlass b/third-party/cutlass
index 62750a2b..6dd13d42 160000
--- a/third-party/cutlass
+++ b/third-party/cutlass
@@ -1 +1 @@
-Subproject commit 62750a2b75c802660e4894434dc55e839f322277
+Subproject commit 6dd13d42784ee5bfa232d2441e6b9a021c5c6290

From 241a40ee6e7569a6ca1f69e95d9c05ed3e62ea3a Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Fri, 8 Aug 2025 10:38:52 +0800
Subject: [PATCH 48/57] Update README.md

---
 README.md | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/README.md b/README.md
index 81cf9243..0da21142 100644
--- a/README.md
+++ b/README.md
@@ -48,11 +48,14 @@
 
 - [2025-01-08]: [📚Split Q + Fully QKV Fine-grained Tiling](#mma-tiling-qkv) has been refactored into 🤖[ffpa-attn](https://github.com/xlite-dev/ffpa-attn.git): 📚FFPA - Yet another Faster Flash Prefill Attention with O(1)🎉SRAM complexity for headdim > 256, **1.8x~3x**🎉faster than SDPA EA: [📈L20 ~1.9x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-l20), [📈 A30 ~1.8x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-a30), [📈3080 ~2.9x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-3080), [📈4090 ~2.1x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-4090).
 
+<!--
 <div align='center'>
   <img src='https://github.com/user-attachments/assets/cba2edce-ac0d-412e-823c-7eea2cc63f83' height="170px" width="229px">
   <img src='https://github.com/user-attachments/assets/447e2937-f7c8-47c8-8550-8c0c71b910e6' height="170px" width="229px">
   <img src='https://github.com/user-attachments/assets/65a8d564-8fa7-4d66-86b9-e238feb86143' height="170px" width="229px">
 </div>
+--> 
+<img width="6412" height="2872" alt="image" src="https://github.com/user-attachments/assets/ed30185b-2e11-4293-832f-43e9003d6ad9" />
 
 - [2024-12-02]: HGEMM MMA kernels has been refactored into 🤖[HGEMM](https://github.com/xlite-dev/HGEMM.git): ⚡️Write HGEMM from scratch using Tensor Cores with WMMA, MMA and CuTe API, achieve peak⚡️ performance.
 

From 0b238d10fc8589169a8c44072cf733c7e8ef62ba Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Fri, 8 Aug 2025 10:41:43 +0800
Subject: [PATCH 49/57] Update README.md

---
 README.md | 4 +++-
 1 file changed, 3 insertions(+), 1 deletion(-)

diff --git a/README.md b/README.md
index 0da21142..bb679fe0 100644
--- a/README.md
+++ b/README.md
@@ -55,7 +55,9 @@
   <img src='https://github.com/user-attachments/assets/65a8d564-8fa7-4d66-86b9-e238feb86143' height="170px" width="229px">
 </div>
 --> 
-<img width="6412" height="2872" alt="image" src="https://github.com/user-attachments/assets/ed30185b-2e11-4293-832f-43e9003d6ad9" />
+<div align='center'>
+<img height="320px" alt="image" src="https://github.com/user-attachments/assets/ed30185b-2e11-4293-832f-43e9003d6ad9" />
+</div>
 
 - [2024-12-02]: HGEMM MMA kernels has been refactored into 🤖[HGEMM](https://github.com/xlite-dev/HGEMM.git): ⚡️Write HGEMM from scratch using Tensor Cores with WMMA, MMA and CuTe API, achieve peak⚡️ performance.
 

From a7d26e6ebacab1c3a880596931439a449514eb35 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Fri, 8 Aug 2025 10:42:44 +0800
Subject: [PATCH 50/57] Update README.md

---
 README.md | 2 --
 1 file changed, 2 deletions(-)

diff --git a/README.md b/README.md
index bb679fe0..5cef9823 100644
--- a/README.md
+++ b/README.md
@@ -38,8 +38,6 @@
 ## 📖 News 🔥🔥
 <div id="news"></div>
 
-- [2025-07-13] **[🤗flux-faster](https://github.com/xlite-dev/flux-faster)** is released! A forked version of [huggingface/flux-fast](https://github.com/huggingface/flux-fast) that **makes flux-fast even faster** with **[cache-dit](https://github.com/vipshop/cache-dit)**, **3.3x** speedup on NVIDIA L20 while still maintaining **high precision**.
-
 - [2025-06-16]: **[🤗cache-dit](https://github.com/vipshop/cache-dit)** is released! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for DiTs (**DBCache**, **DBPrune**, **TaylorSeer**, **FBCache**, **Cache CFG**, etc)🔥. Feel free to take a try!
 
 <div align='center'>

From 1144d0c09e794d7a4138d8dd071ec429609bd1fb Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Fri, 8 Aug 2025 11:24:14 +0800
Subject: [PATCH 51/57] Update README.md

---
 README.md | 7 +++++--
 1 file changed, 5 insertions(+), 2 deletions(-)

diff --git a/README.md b/README.md
index 5cef9823..a438dfe9 100644
--- a/README.md
+++ b/README.md
@@ -39,12 +39,14 @@
 <div id="news"></div>
 
 - [2025-06-16]: **[🤗cache-dit](https://github.com/vipshop/cache-dit)** is released! A **Training-free** and **Easy-to-use** Cache Acceleration Toolbox for DiTs (**DBCache**, **DBPrune**, **TaylorSeer**, **FBCache**, **Cache CFG**, etc)🔥. Feel free to take a try!
+<!--
 
 <div align='center'>
   <img src='https://github.com/user-attachments/assets/a5ec4320-d2f9-4254-888a-170b2d9e3784' height=170px>
 </div>
+--> 
 
-- [2025-01-08]: [📚Split Q + Fully QKV Fine-grained Tiling](#mma-tiling-qkv) has been refactored into 🤖[ffpa-attn](https://github.com/xlite-dev/ffpa-attn.git): 📚FFPA - Yet another Faster Flash Prefill Attention with O(1)🎉SRAM complexity for headdim > 256, **1.8x~3x**🎉faster than SDPA EA: [📈L20 ~1.9x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-l20), [📈 A30 ~1.8x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-a30), [📈3080 ~2.9x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-3080), [📈4090 ~2.1x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-4090).
+- [2025-01-08]: **[🤖ffpa-attn](https://github.com/xlite-dev/ffpa-attn.git)** is released! Yet another Faster Flash Prefill Attention with O(1)🎉SRAM complexity for large headdim, **1.8x~3x↑**🎉 vs SDPA EA: [📈L20 ~1.9x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-l20), [📈A30 ~1.8x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-a30),[📈4090 ~2.1x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-4090).
 
 <!--
 <div align='center'>
@@ -59,12 +61,13 @@
 
 - [2024-12-02]: HGEMM MMA kernels has been refactored into 🤖[HGEMM](https://github.com/xlite-dev/HGEMM.git): ⚡️Write HGEMM from scratch using Tensor Cores with WMMA, MMA and CuTe API, achieve peak⚡️ performance.
 
+<!--
 <div align='center'>
   <img src='https://github.com/user-attachments/assets/71927ac9-72b3-4ce9-b0e2-788b5885bc99' height="170px" width="229px">
   <img src='https://github.com/user-attachments/assets/05ef4f5e-d999-48ea-b58e-782cffb24e85' height="170px" width="229px">
   <img src='https://github.com/user-attachments/assets/9472e970-c083-4b31-9252-3eeecc761078' height="170px" width="229px">
 </div>
-
+-->
 
 ## 📖 Contents
 <div id="contents"></div>

From 0d05779668bf6d55abfdf85468b52d7ac07163aa Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Fri, 8 Aug 2025 11:27:14 +0800
Subject: [PATCH 52/57] Update README.md

---
 README.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/README.md b/README.md
index a438dfe9..b7b4dcdf 100644
--- a/README.md
+++ b/README.md
@@ -59,7 +59,7 @@
 <img height="320px" alt="image" src="https://github.com/user-attachments/assets/ed30185b-2e11-4293-832f-43e9003d6ad9" />
 </div>
 
-- [2024-12-02]: HGEMM MMA kernels has been refactored into 🤖[HGEMM](https://github.com/xlite-dev/HGEMM.git): ⚡️Write HGEMM from scratch using Tensor Cores with WMMA, MMA and CuTe API, achieve peak⚡️ performance.
+- [2024-12-02]: **[⚡️HGEMM](https://github.com/xlite-dev/HGEMM.git)** is released! Write HGEMM from scratch using Tensor Cores with **WMMA, MMA and CuTe** API, achieve peak🎉 performance.
 
 <!--
 <div align='center'>

From e0da7bd02a41fffd942b31b016ddecf4e0c7620a Mon Sep 17 00:00:00 2001
From: leo <92263211+napleon-liu@users.noreply.github.com>
Date: Sat, 9 Aug 2025 09:43:31 +0800
Subject: [PATCH 53/57] fix: softmax.cu spell error (#371)

* softmax.cu spell error

* format

* format
---
 README.md                  |  4 ++--
 kernels/softmax/softmax.cu | 12 ++++++------
 2 files changed, 8 insertions(+), 8 deletions(-)

diff --git a/README.md b/README.md
index b7b4dcdf..f15ea16a 100644
--- a/README.md
+++ b/README.md
@@ -44,7 +44,7 @@
 <div align='center'>
   <img src='https://github.com/user-attachments/assets/a5ec4320-d2f9-4254-888a-170b2d9e3784' height=170px>
 </div>
---> 
+-->
 
 - [2025-01-08]: **[🤖ffpa-attn](https://github.com/xlite-dev/ffpa-attn.git)** is released! Yet another Faster Flash Prefill Attention with O(1)🎉SRAM complexity for large headdim, **1.8x~3x↑**🎉 vs SDPA EA: [📈L20 ~1.9x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-l20), [📈A30 ~1.8x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-a30),[📈4090 ~2.1x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-4090).
 
@@ -54,7 +54,7 @@
   <img src='https://github.com/user-attachments/assets/447e2937-f7c8-47c8-8550-8c0c71b910e6' height="170px" width="229px">
   <img src='https://github.com/user-attachments/assets/65a8d564-8fa7-4d66-86b9-e238feb86143' height="170px" width="229px">
 </div>
---> 
+-->
 <div align='center'>
 <img height="320px" alt="image" src="https://github.com/user-attachments/assets/ed30185b-2e11-4293-832f-43e9003d6ad9" />
 </div>
diff --git a/kernels/softmax/softmax.cu b/kernels/softmax/softmax.cu
index d5c7ed9f..e18761d1 100644
--- a/kernels/softmax/softmax.cu
+++ b/kernels/softmax/softmax.cu
@@ -333,14 +333,14 @@ __global__ void online_safe_softmax_f32_per_token_kernel(const float *x,
   // for softmax)
   int local_tid = threadIdx.x;
   int global_tid = blockIdx.x * NUM_THREADS + threadIdx.x;
-  const int WAPR_NUM = NUM_THREADS / WARP_SIZE;
+  const int WARP_NUM = NUM_THREADS / WARP_SIZE;
   int warp_id = local_tid / WARP_SIZE;
   int lane_id = local_tid % WARP_SIZE;
   MD val;
   val.m = global_tid < N ? x[global_tid] : -FLT_MAX;
   val.d = global_tid < N ? 1.0f : 0.0f;
 
-  __shared__ MD shared[WAPR_NUM];
+  __shared__ MD shared[WARP_NUM];
   MD res = warp_reduce_md_op<WARP_SIZE>(val);
 
   if (lane_id == 0)
@@ -349,7 +349,7 @@ __global__ void online_safe_softmax_f32_per_token_kernel(const float *x,
 
   if (local_tid < WARP_SIZE) {
     MD block_res = shared[local_tid];
-    block_res = warp_reduce_md_op<WAPR_NUM>(block_res);
+    block_res = warp_reduce_md_op<WARP_NUM>(block_res);
     if (local_tid == 0) {
       shared[0] = block_res;
     }
@@ -371,7 +371,7 @@ online_safe_softmax_f32x4_pack_per_token_kernel(float *x, float *y, int N) {
   int local_tid = threadIdx.x;
   int global_tid = (blockIdx.x * NUM_THREADS + local_tid) * 4;
 
-  const int WAPR_NUM = NUM_THREADS / WARP_SIZE;
+  const int WARP_NUM = NUM_THREADS / WARP_SIZE;
   int warp_id = local_tid / WARP_SIZE;
   int lane_id = local_tid % WARP_SIZE;
   // compare local max value
@@ -382,7 +382,7 @@ online_safe_softmax_f32x4_pack_per_token_kernel(float *x, float *y, int N) {
 
   MD local_md = {local_m, local_d};
   MD res = warp_reduce_md_op<WARP_SIZE>(local_md);
-  __shared__ MD shared[WAPR_NUM];
+  __shared__ MD shared[WARP_NUM];
 
   if (lane_id == 0)
     shared[warp_id] = res;
@@ -390,7 +390,7 @@ online_safe_softmax_f32x4_pack_per_token_kernel(float *x, float *y, int N) {
   // do block reduce
   if (local_tid < WARP_SIZE) {
     MD block_res = shared[local_tid];
-    block_res = warp_reduce_md_op<WAPR_NUM>(block_res);
+    block_res = warp_reduce_md_op<WARP_NUM>(block_res);
     if (local_tid == 0)
       shared[0] = block_res;
   }

From 405bad5eb5ab12ad0740800ac3bde4542f7454e6 Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Sat, 9 Aug 2025 12:08:23 +0800
Subject: [PATCH 54/57] feat: taylorseer zhihu blog (#373)

---
 README.md | 1 +
 1 file changed, 1 insertion(+)

diff --git a/README.md b/README.md
index f15ea16a..4a1ccce8 100644
--- a/README.md
+++ b/README.md
@@ -493,6 +493,7 @@ The kernels listed here will guide you through a step-by-step progression, rangi
 
 |📖 类型-标题|📖 作者| 📖 推荐 |
 |:---|:---|:---|
+| [[Diffusion推理]📖Cache加速-TaylorSeer算法简析](https://zhuanlan.zhihu.com/p/1937477466475197176)|@DefTruth|⭐️⭐️⭐|
 | [[Diffusion推理]📖DiT推理加速综述: Caching](https://zhuanlan.zhihu.com/p/711223667)|@DefTruth|⭐️⭐️⭐|
 | [[Triton编程][基础]📖Triton极简入门: Triton Vector Add](https://zhuanlan.zhihu.com/p/1902778199261291694)|@DefTruth|⭐️⭐️⭐|
 | [[Triton编程][基础]📖Triton Fused Softmax Kernel详解: 从Python源码到PTX](https://zhuanlan.zhihu.com/p/1899562146477609112)|@DefTruth|⭐️⭐️⭐|

From fe0c9b5bef53d8f50c58e346d1f3cdb3ff11522d Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Tue, 12 Aug 2025 12:09:05 +0800
Subject: [PATCH 55/57] =?UTF-8?q?cache-dit:=20Qwen-Image=201.5x=E2=86=91?=
 =?UTF-8?q?=20blog=20(#374)?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

---
 README.md | 1 +
 1 file changed, 1 insertion(+)

diff --git a/README.md b/README.md
index 4a1ccce8..f6996973 100644
--- a/README.md
+++ b/README.md
@@ -493,6 +493,7 @@ The kernels listed here will guide you through a step-by-step progression, rangi
 
 |📖 类型-标题|📖 作者| 📖 推荐 |
 |:---|:---|:---|
+| [[Diffusion推理]📖cache-dit: Qwen-Image 1.5x 无损加速](https://zhuanlan.zhihu.com/p/1938547315221705644)|@DefTruth|⭐️⭐️⭐|
 | [[Diffusion推理]📖Cache加速-TaylorSeer算法简析](https://zhuanlan.zhihu.com/p/1937477466475197176)|@DefTruth|⭐️⭐️⭐|
 | [[Diffusion推理]📖DiT推理加速综述: Caching](https://zhuanlan.zhihu.com/p/711223667)|@DefTruth|⭐️⭐️⭐|
 | [[Triton编程][基础]📖Triton极简入门: Triton Vector Add](https://zhuanlan.zhihu.com/p/1902778199261291694)|@DefTruth|⭐️⭐️⭐|

From 0420e93e2a3e3b731a33d022d175bef424bc8acd Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Thu, 14 Aug 2025 08:44:35 +0800
Subject: [PATCH 56/57] Update README.md (#375)

* Update README.md

* Update README.md
---
 README.md | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/README.md b/README.md
index f6996973..24c28295 100644
--- a/README.md
+++ b/README.md
@@ -1,8 +1,7 @@
 <div align="center">
   <p align="center">
     <h2>📚 LeetCUDA: Modern CUDA Learn Notes with PyTorch for Beginners 🐑</h2>
-    <img src='https://github.com/user-attachments/assets/b2578723-b7a7-4d8f-bcd1-5008947b808a' width="360" height="56" >
-    <a href="https://hellogithub.com/repository/98348655a96640ca8ddcbc298edc901d" target="_blank"><img src="https://api.hellogithub.com/v1/widgets/recommend.svg?rid=98348655a96640ca8ddcbc298edc901d&claim_uid=ofSCbzTmdeQk3FD&theme=dark" alt="Featured｜HelloGitHub" style="width: 250px; height: 54px;" width="250" height="54" /></a>
+    <img src='https://github.com/user-attachments/assets/b2578723-b7a7-4d8f-bcd1-5008947b808a' width="700">
   </p>
   <div align='center'>
       <img src=https://cdn.rawgit.com/sindresorhus/awesome/d7305f38d29fed78fa85652e3a63e154dd8e8829/media/badge.svg >
@@ -20,6 +19,7 @@
     <a href="#contribute">🔥🔥 PR Welcome: Add Your Kernel to LeetCUDA! Let's make it Awesome together! 🎉🎉</a> <br>
     <a href=https://github.com/xlite-dev/LeetCUDA/graphs/contributors > <img src=https://opencollective.com/leetcuda/contributors.svg height=40px > </a>
   </p>
+  <a href="https://hellogithub.com/repository/98348655a96640ca8ddcbc298edc901d" target="_blank"><img src="https://api.hellogithub.com/v1/widgets/recommend.svg?rid=98348655a96640ca8ddcbc298edc901d&claim_uid=ofSCbzTmdeQk3FD&theme=dark" alt="Featured｜HelloGitHub" style="width: 250px; height: 54px;" width="250" height="54" /></a> 
 </div>
 
 ## ©️Citations🎉🎉

From fb78f162b93db026b2bcd3f85df8c2fe3a5a8eab Mon Sep 17 00:00:00 2001
From: DefTruth <31974251+DefTruth@users.noreply.github.com>
Date: Thu, 14 Aug 2025 19:42:01 +0800
Subject: [PATCH 57/57] Update README.md

---
 README.md | 26 ++++++++++++++++++++++++++
 1 file changed, 26 insertions(+)

diff --git a/README.md b/README.md
index 24c28295..dc726d77 100644
--- a/README.md
+++ b/README.md
@@ -1,3 +1,4 @@
+<!--
 <div align="center">
   <p align="center">
     <h2>📚 LeetCUDA: Modern CUDA Learn Notes with PyTorch for Beginners 🐑</h2>
@@ -21,6 +22,31 @@
   </p>
   <a href="https://hellogithub.com/repository/98348655a96640ca8ddcbc298edc901d" target="_blank"><img src="https://api.hellogithub.com/v1/widgets/recommend.svg?rid=98348655a96640ca8ddcbc298edc901d&claim_uid=ofSCbzTmdeQk3FD&theme=dark" alt="Featured｜HelloGitHub" style="width: 250px; height: 54px;" width="250" height="54" /></a> 
 </div>
+-->
+
+<div align="center">
+  <p align="center">
+    <h2>📚 LeetCUDA: Modern CUDA Learn Notes with PyTorch for Beginners 🐑</h2>
+    <img src='https://github.com/user-attachments/assets/b2578723-b7a7-4d8f-bcd1-5008947b808a' width="360" height="56" >
+    <a href="https://hellogithub.com/repository/98348655a96640ca8ddcbc298edc901d" target="_blank"><img src="https://api.hellogithub.com/v1/widgets/recommend.svg?rid=98348655a96640ca8ddcbc298edc901d&claim_uid=ofSCbzTmdeQk3FD&theme=dark" alt="Featured｜HelloGitHub" style="width: 250px; height: 54px;" width="250" height="54" /></a>
+  </p>
+  <div align='center'>
+      <img src=https://cdn.rawgit.com/sindresorhus/awesome/d7305f38d29fed78fa85652e3a63e154dd8e8829/media/badge.svg >
+      <img src=https://img.shields.io/badge/Language-CUDA-brightgreen.svg >
+      <img src=https://img.shields.io/github/forks/xlite-dev/LeetCUDA.svg?style=dark >
+      <img src=https://img.shields.io/github/stars/xlite-dev/LeetCUDA.svg?style=dark >
+      <img src=https://img.shields.io/badge/License-GPLv3.0-turquoise.svg >
+  </div>
+</div>
+
+📚 **LeetCUDA**: It includes **Tensor/CUDA Cores, TF32/F16/BF16/F8**, [📖200+ CUDA Kernels🔥](#cuda-kernel) with PyTorch, [📖100+ LLM/CUDA🔥](#my-blogs-part-1) blogs, [📖HGEMM⚡️](./kernels/hgemm) which can achieve `98%~100%` TFLOPS of **cuBLAS**, and [📖flash-attn⚡️](./kernels/flash-attn) using Tensor Cores with pure MMA PTX. ♥️ Please consider to leave a ⭐️ Star to support me, my bro ~ ♥️
+
+<div align="center">
+  <p align="center">
+    <a href="#contribute">🔥🔥 PR Welcome: Add Your Kernel to LeetCUDA! Let's make it Awesome together! 🎉🎉</a> <br>
+    <a href=https://github.com/xlite-dev/LeetCUDA/graphs/contributors > <img src=https://opencollective.com/leetcuda/contributors.svg height=40px > </a>
+  </p>
+</div>
 
 ## ©️Citations🎉🎉
 
