[HGEMM] Release toy-hgemm library 0.1.0 (#146)

* refactor * refactor * refactor * refactor * refactor * refactor * refactor * refactor * refactor * refactor * refactor * Update README.md * Update README.md * Update README.md * Update README.md * Update README.md * Update README.md
DefTruth · Nov 22, 2024 · 56e2fe9 · 56e2fe9
1 parent 6ea2eb9
commit 56e2fe9
Show file tree

Hide file tree

Showing 197 changed files with 222 additions and 202 deletions.
diff --git a/.github/workflows/issue.yml b/.github/workflows/issue.yml
@@ -19,4 +19,4 @@ jobs:
         close-issue-message: "This issue was closed because it has been inactive for 7 days since being marked as stale."
         days-before-pr-stale: -1
         days-before-pr-close: -1
-        repo-token: ${{ secrets.GITHUB_TOKEN }}
+        repo-token: ${{ secrets.GITHUB_TOKEN }}
diff --git a/.gitignore b/.gitignore
@@ -16,3 +16,5 @@ __pycache__
 *.ncu*
 *.sqlite*
 *.engine
+*.bin 
+outupt
diff --git a/.gitmodules b/.gitmodules
@@ -1,3 +1,4 @@
 [submodule "third-party/cutlass"]
 	path = third-party/cutlass
 	url = https://github.com/NVIDIA/cutlass.git
+
diff --git a/LICENSE b/LICENSE
@@ -671,4 +671,4 @@ into proprietary programs.  If your program is a subroutine library, you
 may consider it more useful to permit linking proprietary applications with
 the library.  If this is what you want to do, use the GNU Lesser General
 Public License instead of this License.  But first, please read
-<https://www.gnu.org/licenses/why-not-lgpl.html>.
+<https://www.gnu.org/licenses/why-not-lgpl.html>.
diff --git a/README.md b/README.md
diff --git a/hgemm/tools/clear.sh b/hgemm/tools/clear.sh
diff --git a/cuda-slides/.gitignore → kernels/cutlass/.gitignore b/cuda-slides/.gitignore → kernels/cutlass/.gitignore
diff --git a/cutlass/cute/.gitignore → kernels/cutlass/cute/.gitignore b/cutlass/cute/.gitignore → kernels/cutlass/cute/.gitignore
diff --git a/cutlass/cute/vector_add.cu → kernels/cutlass/cute/vector_add.cu b/cutlass/cute/vector_add.cu → kernels/cutlass/cute/vector_add.cu
diff --git a/cutlass/cutlass-3.x/.gitignore → kernels/cutlass/cutlass-3.x/.gitignore b/cutlass/cutlass-3.x/.gitignore → kernels/cutlass/cutlass-3.x/.gitignore
diff --git a/dot-product/.gitignore → kernels/dot-product/.gitignore b/dot-product/.gitignore → kernels/dot-product/.gitignore
diff --git a/dot-product/README.md → kernels/dot-product/README.md b/dot-product/README.md → kernels/dot-product/README.md
diff --git a/dot-product/dot_product.cu → kernels/dot-product/dot_product.cu b/dot-product/dot_product.cu → kernels/dot-product/dot_product.cu
diff --git a/dot-product/dot_product.py → kernels/dot-product/dot_product.py b/dot-product/dot_product.py → kernels/dot-product/dot_product.py
diff --git a/elementwise/.gitignore → kernels/elementwise/.gitignore b/elementwise/.gitignore → kernels/elementwise/.gitignore
diff --git a/elementwise/README.md → kernels/elementwise/README.md b/elementwise/README.md → kernels/elementwise/README.md
diff --git a/elementwise/elementwise.cu → kernels/elementwise/elementwise.cu b/elementwise/elementwise.cu → kernels/elementwise/elementwise.cu
diff --git a/elementwise/elementwise.py → kernels/elementwise/elementwise.py b/elementwise/elementwise.py → kernels/elementwise/elementwise.py
diff --git a/embedding/.gitignore → kernels/embedding/.gitignore b/embedding/.gitignore → kernels/embedding/.gitignore
diff --git a/embedding/README.md → kernels/embedding/README.md b/embedding/README.md → kernels/embedding/README.md
diff --git a/embedding/embedding.cu → kernels/embedding/embedding.cu b/embedding/embedding.cu → kernels/embedding/embedding.cu
diff --git a/embedding/embedding.py → kernels/embedding/embedding.py b/embedding/embedding.py → kernels/embedding/embedding.py
diff --git a/cutlass/.gitignore → kernels/flash-attn/.gitignore b/cutlass/.gitignore → kernels/flash-attn/.gitignore
diff --git a/flash-attn/README.md → kernels/flash-attn/README.md b/flash-attn/README.md → kernels/flash-attn/README.md
diff --git a/flash-attn/flash_attn.cc → kernels/flash-attn/flash_attn.cc b/flash-attn/flash_attn.cc → kernels/flash-attn/flash_attn.cc
diff --git a/flash-attn/flash_attn.cu → kernels/flash-attn/flash_attn.cu b/flash-attn/flash_attn.cu → kernels/flash-attn/flash_attn.cu
diff --git a/flash-attn/flash_attn.py → kernels/flash-attn/flash_attn.py b/flash-attn/flash_attn.py → kernels/flash-attn/flash_attn.py
diff --git a/flash-attn/flash_attn_mma.cu → kernels/flash-attn/flash_attn_mma.cu b/flash-attn/flash_attn_mma.cu → kernels/flash-attn/flash_attn_mma.cu
diff --git a/gelu/.gitignore → kernels/gelu/.gitignore b/gelu/.gitignore → kernels/gelu/.gitignore
diff --git a/gelu/README.md → kernels/gelu/README.md b/gelu/README.md → kernels/gelu/README.md
diff --git a/gelu/gelu.cu → kernels/gelu/gelu.cu b/gelu/gelu.cu → kernels/gelu/gelu.cu
diff --git a/gelu/gelu.py → kernels/gelu/gelu.py b/gelu/gelu.py → kernels/gelu/gelu.py
diff --git a/hgemm/.gitignore → kernels/hgemm/.gitignore b/hgemm/.gitignore → kernels/hgemm/.gitignore
diff --git a/hgemm/README.md → kernels/hgemm/README.md b/hgemm/README.md → kernels/hgemm/README.md
@@ -1,56 +1,65 @@
-# 🔥🔥Toy-HGEMM Library: Achieve the performance of cuBLAS
+## 🔥🔥Toy-HGEMM Library: Achieve the performance of cuBLAS
 
 |CUDA Cores|Sliced K(Loop over K)|Tile Block|Tile Thread|
 |:---:|:---:|:---:|:---:|
 |✔️|✔️|✔️|✔️|
-|**WMMA(m16n16k16)**|**MMA(m16n8k16)**|**Pack LDST(128 bits)**|**SMEM Padding**|
+|WMMA(m16n16k16)|MMA(m16n8k16)|Pack LDST(128 bits)|SMEM Padding|
 |✔️|✔️|✔️|✔️|
-|**Copy Async**|**Tile MMA(More Threads)**|**Tile Warp(More Values)**|**Multi Stages**|  
+|Copy Async|Tile MMA(More Threads)|Tile Warp(More Values)|Multi Stages|  
 |✔️|✔️|✔️|✔️|
-|**Reg Double Buffers**|**Block Swizzle**|**Warp Swizzle**|**Collective Store(Reg Reuse&Warp Shfl)**|
+|Reg Double Buffers|Block Swizzle|Warp Swizzle|Collective Store(Warp Shfl)|
 |✔️|✔️|✔️|✔️|
-|**Row Major(NN)**|**Col Major(TN)**|**SGEMM TF32**|**SMEM Swizzle(CuTe)**|
+|Row Major(NN)|Col Major(TN)|SGEMM TF32|SMEM Swizzle(CuTe)|
 |✔️|✔️|✔️|✔️|
 
-<details>
-<summary> 🔑️ 点击查看所有支持的HGEMM Kernels! </summary>  
-
-- [X] hgemm_sliced_k_f16_kernel 
-- [X] hgemm_t_8x8_sliced_k_f16x4_kernel(unpack)
-- [X] hgemm_t_8x8_sliced_k_f16x4_pack_kernel(pack 16x4)
-- [X] hgemm_t_8x8_sliced_k_f16x4_bcf_kernel(bank conflicts reduce)
-- [X] hgemm_t_8x8_sliced_k_f16x4_pack_bcf_kernel(bank conflicts reduce, pack)
-- [X] hgemm_t_8x8_sliced_k_f16x8_pack_bcf_kernel(bank conflicts reduce, pack)
-- [X] hgemm_t_8x8_sliced_k_f16x8_pack_bcf_dbuf_kernel(bank conflicts reduce, pack, double buffers)
-- [X] hgemm_t_8x8_sliced_k16/32_f16x8_pack_bcf_dbuf_kernel(pack, double buffers)
-- [X] hgemm_t_8x8_sliced_k16/32_f16x8_pack_bcf_dbuf_async_kernel(pack, double buffers, copy async)
-- [X] hgemm_wmma_m16n16k16_naive(WMMA) 
-- [X] hgemm_wmma_m16n16k16_mma4x2(WMMA, Tile MMA) 
-- [X] hgemm_wmma_m16n16k16_mma4x2_warp2x4(TWMMA, Tile MMA/Warp, pack) 
-- [X] hgemm_wmma_m16n16k16_mma4x2_warp2x4_async(WMMA, Tile MMA/Warp, Copy Async) 
-- [X] hgemm_wmma_m16n16k16_mma4x2_warp2x4_async_offset(WMMA, Tile MMA/Warp, Copy Async, Pad)
-- [X] hgemm_wmma_m16n16k16_mma4x2_warp2x4_dbuf_async(WMMA, Tile MMA/Warp, Copy Async, Double Buffers, Pad)  
-- [X] hgemm_wmma_m16n16k16_mma4x2_warp2x4_stages(WMMA, Tile MMA/Warp, Copy Async, Stages, Pad, Block swizzle) 
-- [X] hgemm_wmma_m16n16k16_mma4x2_warp4x4_stages(WMMA, Tile MMA/Warp, Copy Async, Stages, Pad, Block swizzle)
-- [X] hgemm_wmma_m16n16k16_mma4x4_warp4x4_stages(WMMA, Tile MMA/Warp, Copy Async, Stages, Pad, Block swizzle) 
-- [X] hgemm_wmma_m32n8k16_mma2x4_warp2x4_dbuf_async(WMMA, Tile MMA/Warp, Copy Async, Double Buffers, Pad)
-- [X] hgemm_mma_m16n8k16_naive(MMA)
-- [X] hgemm_mma_m16n8k16_mma2x4_warp4x4(MMA, Tile MMA/Warp, pack)
-- [X] hgemm_mma_m16n8k16_mma2x4_warp4x4_stages(MMA, Tile MMA/Warp, Copy Async, Stages, Pad, Block swizzle)
-- [X] hgemm_mma_m16n8k16_mma2x4_warp4x4x2_stages(MMA, Tile MMA/Warp, Copy Async, Stages, Pad, Block swizzle, Warp swizzle, Reg Double Buffers, Collective Store with Reg Reuse & Warp Shuffle) 
-- [X] hgemm_mma_stages_block_swizzle_tn_cute(MMA, Tile MMA/Warp, Copy Async, Stages, Block Swizzle, SMEM Swizzle, Collective Store with SMEM) 
-- [X] PyTorch bindings
 
-</details>
+## 📖 HGEMM CUDA Kernels in Toy-HGEMM Library 🎉🎉 
+
+```C++  
+void hgemm_naive_f16(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_sliced_k_f16(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_t_8x8_sliced_k_f16x4(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_t_8x8_sliced_k_f16x4_pack(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_t_8x8_sliced_k_f16x4_bcf(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_t_8x8_sliced_k_f16x4_pack_bcf(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_t_8x8_sliced_k_f16x8_pack_bcf(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_t_8x8_sliced_k_f16x8_pack_bcf_dbuf(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_t_8x8_sliced_k16_f16x8_pack_dbuf(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_t_8x8_sliced_k16_f16x8_pack_dbuf_async(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_t_8x8_sliced_k32_f16x8_pack_dbuf(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_t_8x8_sliced_k32_f16x8_pack_dbuf_async(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_t_16x8_sliced_k32_f16x8_pack_dbuf(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_t_16x8_sliced_k32_f16x8_pack_dbuf_async(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_cublas_tensor_op_nn(torch::Tensor a, torch::Tensor b, torch::Tensor c); 
+void hgemm_cublas_tensor_op_tn(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_wmma_m16n16k16_naive(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_wmma_m16n16k16_mma4x2(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_wmma_m16n16k16_mma4x2_warp2x4(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_wmma_m16n16k16_mma4x2_warp2x4_dbuf_async(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_wmma_m32n8k16_mma2x4_warp2x4_dbuf_async(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_wmma_m16n16k16_mma4x2_warp2x4_stages(torch::Tensor a, torch::Tensor b, torch::Tensor c, int stages, bool swizzle, int swizzle_stride);
+void hgemm_wmma_m16n16k16_mma4x2_warp2x4_stages_dsmem(torch::Tensor a, torch::Tensor b, torch::Tensor c, int stages, bool swizzle, int swizzle_stride);
+void hgemm_wmma_m16n16k16_mma4x2_warp4x4_stages_dsmem(torch::Tensor a, torch::Tensor b, torch::Tensor c, int stages, bool swizzle, int swizzle_stride);                                                        
+void hgemm_wmma_m16n16k16_mma4x4_warp4x4_stages_dsmem(torch::Tensor a, torch::Tensor b, torch::Tensor c, int stages, bool swizzle, int swizzle_stride);
+void hgemm_mma_m16n8k16_naive(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_mma_m16n8k16_mma2x4_warp4x4(torch::Tensor a, torch::Tensor b, torch::Tensor c);
+void hgemm_mma_m16n8k16_mma2x4_warp4x4_stages(torch::Tensor a, torch::Tensor b, torch::Tensor c, int stages, bool swizzle, int swizzle_stride);
+void hgemm_mma_m16n8k16_mma2x4_warp4x4_stages_dsmem(torch::Tensor a, torch::Tensor b, torch::Tensor c, int stages, bool swizzle, int swizzle_stride);
+void hgemm_mma_m16n8k16_mma2x4_warp4x4x2_stages_dsmem(torch::Tensor a, torch::Tensor b, torch::Tensor c, int stages, bool swizzle, int swizzle_stride);
+void hgemm_mma_m16n8k16_mma2x4_warp4x4x2_stages_dsmem_x4(torch::Tensor a, torch::Tensor b, torch::Tensor c, int stages, bool swizzle, int swizzle_stride);
+void hgemm_mma_m16n8k16_mma2x4_warp4x4x2_stages_dsmem_rr(torch::Tensor a, torch::Tensor b, torch::Tensor c, int stages, bool swizzle, int swizzle_stride);
+void hgemm_mma_m16n8k16_mma2x4_warp4x4_stages_dsmem_tn(torch::Tensor a, torch::Tensor b, torch::Tensor c, int stages, bool swizzle, int swizzle_stride);
+void hgemm_mma_stages_tn_cute(torch::Tensor a, torch::Tensor b, torch::Tensor c, int stages, bool swizzle, int swizzle_stride);
+```
 
-## 安装
-本仓库实现的HGEMM CUDA kernels可以作为一个python库toy-hgemm使用，安装命令如下。（可选）
+## 📖 安装
+本仓库实现的HGEMM可以作为一个python库使用（可选）
 ```bash
-git submodule update --init --recursive --force
-bash tools/install.sh # pip uninstall toy-hgemm 卸载
+git submodule update --init --recursive --force # 更新cutlass, 必须
+python3 setup.py bdist_wheel && cd dist && python3 -m pip install *.whl # pip uninstall toy-hgemm -y 卸载
 ```
 
-## 测试命令
+## 📖 测试
 
 **CUTLASS**: 更新CUTLASS依赖库
 ```bash
@@ -114,7 +123,7 @@ M N K =  16128  16128  16128, Time =   0.07319142   0.07320709   0.07326925 s, A
 M N K =  16384  16384  16384, Time =   0.07668429   0.07669371   0.07670784 s, AVG Performance =   114.6912 Tflops
 ```
 
-## 目前性能  
+## 📖 目前性能  
 
 ### NVIDIA L20  
 
@@ -132,8 +141,6 @@ M N K =  16384  16384  16384, Time =   0.07668429   0.07669371   0.07670784 s, A
 
 ![NVIDIA_L20_NN+TN+v2](https://github.com/user-attachments/assets/71927ac9-72b3-4ce9-b0e2-788b5885bc99)
 
-- WMMA: Up to 113.76 TFLOPS, 113.83/119.5=95.25% TFLOPS utilization, 113.83/116.25=97.91% cuBLAS performance.
-- MMA: Up to 115.12 TFLOPS, 115.12/119.5=96.33% TFLOPS utilization, 115.12/116.25=99.03% cuBLAS performance.
 
 全量MNK测试命令（提示: 每个MNK单独测试的性能数据更准确）
 ```bash
@@ -166,7 +173,7 @@ python3 hgemm.py --wmma-all --plot
 ```
 
 
-## 性能优化笔记
+## 📖 性能优化笔记
 
 ### PyTorch HGEMM Profile
 

diff --git a/...IDIA_GeForce_RTX_3080_Laptop_GPU_WSL2.png → ...IDIA_GeForce_RTX_3080_Laptop_GPU_WSL2.png b/...IDIA_GeForce_RTX_3080_Laptop_GPU_WSL2.png → ...IDIA_GeForce_RTX_3080_Laptop_GPU_WSL2.png
diff --git a/hgemm/bench/NVIDIA_GeForce_RTX_4090.png → ...s/hgemm/bench/NVIDIA_GeForce_RTX_4090.png b/hgemm/bench/NVIDIA_GeForce_RTX_4090.png → ...s/hgemm/bench/NVIDIA_GeForce_RTX_4090.png
diff --git a/hgemm/bench/NVIDIA_L20.png → kernels/hgemm/bench/NVIDIA_L20.png b/hgemm/bench/NVIDIA_L20.png → kernels/hgemm/bench/NVIDIA_L20.png
diff --git a/hgemm/bench/prof.py → kernels/hgemm/bench/prof.py b/hgemm/bench/prof.py → kernels/hgemm/bench/prof.py
diff --git a/hgemm/cublas/hgemm_cublas.cu → kernels/hgemm/cublas/hgemm_cublas.cu b/hgemm/cublas/hgemm_cublas.cu → kernels/hgemm/cublas/hgemm_cublas.cu
diff --git a/hgemm/cutlass/hgemm_mma_stage_tn_cute.cu → .../hgemm/cutlass/hgemm_mma_stage_tn_cute.cu b/hgemm/cutlass/hgemm_mma_stage_tn_cute.cu → .../hgemm/cutlass/hgemm_mma_stage_tn_cute.cu
diff --git a/hgemm/hgemm.py → kernels/hgemm/hgemm.py b/hgemm/hgemm.py → kernels/hgemm/hgemm.py
diff --git a/hgemm/makefile → kernels/hgemm/makefile b/hgemm/makefile → kernels/hgemm/makefile
@@ -1,4 +1,4 @@
-INCLUDE_DIRS=-I ../ -I ./utils -I ../third-party/cutlass/include -I ../third-party/cutlass/tools/util/include
+INCLUDE_DIRS=-I ./utils -I ../../third-party/cutlass/include -I ../../third-party/cutlass/tools/util/include
 default:
 	nvcc cutlass/hgemm_mma_stage_tn_cute.cu -o hgemm_cute.bin -O2 -arch=sm_89 -std=c++17 $(INCLUDE_DIRS) --expt-relaxed-constexpr -lcublas
 	nvcc cublas/hgemm_cublas.cu -o hgemm_cublas.bin -O2 -arch=sm_89 -std=c++17 $(INCLUDE_DIRS) --expt-relaxed-constexpr -lcublas

diff --git a/hgemm/mma/hgemm_mma.cu → kernels/hgemm/mma/hgemm_mma.cu b/hgemm/mma/hgemm_mma.cu → kernels/hgemm/mma/hgemm_mma.cu
diff --git a/hgemm/mma/hgemm_mma_stage.cu → kernels/hgemm/mma/hgemm_mma_stage.cu b/hgemm/mma/hgemm_mma_stage.cu → kernels/hgemm/mma/hgemm_mma_stage.cu
diff --git a/hgemm/mma/hgemm_mma_stage_tn.cu → kernels/hgemm/mma/hgemm_mma_stage_tn.cu b/hgemm/mma/hgemm_mma_stage_tn.cu → kernels/hgemm/mma/hgemm_mma_stage_tn.cu
diff --git a/hgemm/naive/hgemm.cu → kernels/hgemm/naive/hgemm.cu b/hgemm/naive/hgemm.cu → kernels/hgemm/naive/hgemm.cu
diff --git a/hgemm/naive/hgemm_async.cu → kernels/hgemm/naive/hgemm_async.cu b/hgemm/naive/hgemm_async.cu → kernels/hgemm/naive/hgemm_async.cu
diff --git a/hgemm/pybind/hgemm.cc → kernels/hgemm/pybind/hgemm.cc b/hgemm/pybind/hgemm.cc → kernels/hgemm/pybind/hgemm.cc
diff --git a/hgemm/setup.py → kernels/hgemm/setup.py b/hgemm/setup.py → kernels/hgemm/setup.py
@@ -10,7 +10,7 @@
 )
 from tools.utils import (get_build_sources, get_build_cuda_cflags)
 
-# package name managed by pip, which can be remove by `pip uninstall tiny_pkg`
+# package name managed by pip, which can be remove by `pip uninstall toy-hgemm`
 PACKAGE_NAME = "toy-hgemm"
 
 ext_modules = []

diff --git a/kernels/hgemm/tools/clear.sh b/kernels/hgemm/tools/clear.sh
@@ -0,0 +1,5 @@
+set -x
+
+rm -rf __pycache__ build dist toy_hgemm.egg-info *.bin
+
+set +x
diff --git a/hgemm/tools/install.sh → kernels/hgemm/tools/install.sh b/hgemm/tools/install.sh → kernels/hgemm/tools/install.sh
@@ -1,4 +1,8 @@
+set -x
+
+git submodule update --init --recursive --force
 python3 -m pip uninstall toy-hgemm -y
 python3 setup.py bdist_wheel && cd dist && python3 -m pip install *.whl && cd -
 rm -rf toy_hgemm.egg-info __pycache__
 
+set +x
diff --git a/hgemm/tools/utils.py → kernels/hgemm/tools/utils.py b/hgemm/tools/utils.py → kernels/hgemm/tools/utils.py
@@ -31,8 +31,8 @@ def get_build_sources():
 
 
 def get_project_dir():
-    return os.path.dirname(
-        os.path.dirname(os.path.dirname(os.path.abspath(__file__))))
+    return os.path.dirname(os.path.dirname(
+        os.path.dirname(os.path.dirname(os.path.abspath(__file__)))))
 
 
 def get_build_cuda_cflags(build_pkg: bool = False):
@@ -68,22 +68,22 @@ def get_build_cuda_cflags(build_pkg: bool = False):
       extra_cuda_cflags.append("-Xptxas -v")
     else:
       extra_cuda_cflags.append("--ptxas-options=-v")
-      extra_cuda_cflags.append("--ptxas-options=-O2")
+      extra_cuda_cflags.append("--ptxas-options=-O3")
     # extra cuda flags for cute hgemm
     project_dir = get_project_dir()
     extra_cuda_cflags.append('-DNO_MMA_HGEMM_BIN')
     extra_cuda_cflags.append('-DNO_WMMA_HGEMM_BIN')
     extra_cuda_cflags.append('-DNO_CUTE_HGEMM_BIN')
     extra_cuda_cflags.append('-DNO_CUBLAS_HGEMM_BIN')
     # add cutlass headers and link cublas.
-    extra_cuda_cflags.append(f'-I {project_dir}')
-    extra_cuda_cflags.append(f'-I {project_dir}/utils')
-    extra_cuda_cflags.append(f'-I {project_dir}/naive')
-    extra_cuda_cflags.append(f'-I {project_dir}/wmma')
-    extra_cuda_cflags.append(f'-I {project_dir}/mma')
-    extra_cuda_cflags.append(f'-I {project_dir}/cutlass')
-    extra_cuda_cflags.append(f'-I {project_dir}/cublas')
-    extra_cuda_cflags.append(f'-I {project_dir}/pybind')
+    extra_cuda_cflags.append(f'-I {project_dir}/kernels/hgemm')
+    extra_cuda_cflags.append(f'-I {project_dir}/kernels/hgemm/utils')
+    extra_cuda_cflags.append(f'-I {project_dir}/kernels/hgemm/naive')
+    extra_cuda_cflags.append(f'-I {project_dir}/kernels/hgemm/wmma')
+    extra_cuda_cflags.append(f'-I {project_dir}/kernels/hgemm/mma')
+    extra_cuda_cflags.append(f'-I {project_dir}/kernels/hgemm/cutlass')
+    extra_cuda_cflags.append(f'-I {project_dir}/kernels/hgemm/cublas')
+    extra_cuda_cflags.append(f'-I {project_dir}/kernels/hgemm/pybind')
     extra_cuda_cflags.append(f'-I {project_dir}/third-party/cutlass/include')
     extra_cuda_cflags.append(f'-I {project_dir}/third-party/cutlass/tools/util/include')
     extra_cuda_cflags.append('-lcublas')

diff --git a/hgemm/utils/utils.h → kernels/hgemm/utils/utils.h b/hgemm/utils/utils.h → kernels/hgemm/utils/utils.h
diff --git a/hgemm/wmma/hgemm_wmma.cu → kernels/hgemm/wmma/hgemm_wmma.cu b/hgemm/wmma/hgemm_wmma.cu → kernels/hgemm/wmma/hgemm_wmma.cu
diff --git a/hgemm/wmma/hgemm_wmma_stage.cu → kernels/hgemm/wmma/hgemm_wmma_stage.cu b/hgemm/wmma/hgemm_wmma_stage.cu → kernels/hgemm/wmma/hgemm_wmma_stage.cu
diff --git a/hgemv/.gitignore → kernels/hgemv/.gitignore b/hgemv/.gitignore → kernels/hgemv/.gitignore
diff --git a/hgemv/README.md → kernels/hgemv/README.md b/hgemv/README.md → kernels/hgemv/README.md
diff --git a/hgemv/hgemv.cu → kernels/hgemv/hgemv.cu b/hgemv/hgemv.cu → kernels/hgemv/hgemv.cu
diff --git a/hgemv/hgemv.py → kernels/hgemv/hgemv.py b/hgemv/hgemv.py → kernels/hgemv/hgemv.py
diff --git a/histogram/.gitignore → kernels/histogram/.gitignore b/histogram/.gitignore → kernels/histogram/.gitignore
diff --git a/histogram/README.md → kernels/histogram/README.md b/histogram/README.md → kernels/histogram/README.md
diff --git a/histogram/histogram.cu → kernels/histogram/histogram.cu b/histogram/histogram.cu → kernels/histogram/histogram.cu
diff --git a/histogram/histogram.py → kernels/histogram/histogram.py b/histogram/histogram.py → kernels/histogram/histogram.py
diff --git a/layer-norm/.gitignore → kernels/layer-norm/.gitignore b/layer-norm/.gitignore → kernels/layer-norm/.gitignore
diff --git a/layer-norm/README.md → kernels/layer-norm/README.md b/layer-norm/README.md → kernels/layer-norm/README.md
diff --git a/layer-norm/layer_norm.cu → kernels/layer-norm/layer_norm.cu b/layer-norm/layer_norm.cu → kernels/layer-norm/layer_norm.cu
diff --git a/layer-norm/layer_norm.py → kernels/layer-norm/layer_norm.py b/layer-norm/layer_norm.py → kernels/layer-norm/layer_norm.py
diff --git a/mat-transpose/.gitignore → kernels/mat-transpose/.gitignore b/mat-transpose/.gitignore → kernels/mat-transpose/.gitignore
diff --git a/mat-transpose/README.md → kernels/mat-transpose/README.md b/mat-transpose/README.md → kernels/mat-transpose/README.md
diff --git a/mat-transpose/mat_transpose.cu → kernels/mat-transpose/mat_transpose.cu b/mat-transpose/mat_transpose.cu → kernels/mat-transpose/mat_transpose.cu
diff --git a/mat-transpose/mat_transpose.py → kernels/mat-transpose/mat_transpose.py b/mat-transpose/mat_transpose.py → kernels/mat-transpose/mat_transpose.py
diff --git a/flash-attn/.gitignore → kernels/nms/.gitignore b/flash-attn/.gitignore → kernels/nms/.gitignore
diff --git a/nms/README.md → kernels/nms/README.md b/nms/README.md → kernels/nms/README.md
diff --git a/nms/nms.cc → kernels/nms/nms.cc b/nms/nms.cc → kernels/nms/nms.cc
diff --git a/nms/nms.cu → kernels/nms/nms.cu b/nms/nms.cu → kernels/nms/nms.cu
diff --git a/nms/nms.py → kernels/nms/nms.py b/nms/nms.py → kernels/nms/nms.py
diff --git a/notes-v1.cu → kernels/notes-v1.cu b/notes-v1.cu → kernels/notes-v1.cu
diff --git a/nvidia-nsight/.gitignore → kernels/nvidia-nsight/.gitignore b/nvidia-nsight/.gitignore → kernels/nvidia-nsight/.gitignore
diff --git a/nvidia-nsight/README.md → kernels/nvidia-nsight/README.md b/nvidia-nsight/README.md → kernels/nvidia-nsight/README.md
diff --git a/nvidia-nsight/elementwise.cu → kernels/nvidia-nsight/elementwise.cu b/nvidia-nsight/elementwise.cu → kernels/nvidia-nsight/elementwise.cu
diff --git a/nvidia-nsight/relu.cu → kernels/nvidia-nsight/relu.cu b/nvidia-nsight/relu.cu → kernels/nvidia-nsight/relu.cu
diff --git a/nms/.gitignore → kernels/openai-triton/.gitignore b/nms/.gitignore → kernels/openai-triton/.gitignore
diff --git a/openai-triton/README.md → kernels/openai-triton/README.md b/openai-triton/README.md → kernels/openai-triton/README.md
diff --git a/openai-triton/flash_attn_v2_fwd.py → kernels/openai-triton/flash_attn_v2_fwd.py b/openai-triton/flash_attn_v2_fwd.py → kernels/openai-triton/flash_attn_v2_fwd.py
diff --git a/openai-triton/prefix_prefill.py → kernels/openai-triton/prefix_prefill.py b/openai-triton/prefix_prefill.py → kernels/openai-triton/prefix_prefill.py
diff --git a/openai-triton/prefix_prefill_alibi.py → ...els/openai-triton/prefix_prefill_alibi.py b/openai-triton/prefix_prefill_alibi.py → ...els/openai-triton/prefix_prefill_alibi.py
diff --git a/openai-triton/requirements.txt → kernels/openai-triton/requirements.txt b/openai-triton/requirements.txt → kernels/openai-triton/requirements.txt
diff --git a/openai-triton/test_flash_attn_v2_fwd.py → ...s/openai-triton/test_flash_attn_v2_fwd.py b/openai-triton/test_flash_attn_v2_fwd.py → ...s/openai-triton/test_flash_attn_v2_fwd.py
diff --git a/openai-triton/test_prefix_prefill.py → kernels/openai-triton/test_prefix_prefill.py b/openai-triton/test_prefix_prefill.py → kernels/openai-triton/test_prefix_prefill.py
diff --git a/openai-triton/test_prefix_prefill_alibi.py → ...penai-triton/test_prefix_prefill_alibi.py b/openai-triton/test_prefix_prefill_alibi.py → ...penai-triton/test_prefix_prefill_alibi.py
diff --git a/reduce/.gitignore → kernels/reduce/.gitignore b/reduce/.gitignore → kernels/reduce/.gitignore
diff --git a/reduce/README.md → kernels/reduce/README.md b/reduce/README.md → kernels/reduce/README.md
diff --git a/reduce/block_all_reduce.cu → kernels/reduce/block_all_reduce.cu b/reduce/block_all_reduce.cu → kernels/reduce/block_all_reduce.cu
diff --git a/reduce/block_all_reduce.py → kernels/reduce/block_all_reduce.py b/reduce/block_all_reduce.py → kernels/reduce/block_all_reduce.py
diff --git a/relu/.gitignore → kernels/relu/.gitignore b/relu/.gitignore → kernels/relu/.gitignore
diff --git a/relu/README.md → kernels/relu/README.md b/relu/README.md → kernels/relu/README.md
diff --git a/relu/relu.cu → kernels/relu/relu.cu b/relu/relu.cu → kernels/relu/relu.cu
diff --git a/relu/relu.py → kernels/relu/relu.py b/relu/relu.py → kernels/relu/relu.py
diff --git a/rms-norm/.gitignore → kernels/rms-norm/.gitignore b/rms-norm/.gitignore → kernels/rms-norm/.gitignore
diff --git a/rms-norm/README.md → kernels/rms-norm/README.md b/rms-norm/README.md → kernels/rms-norm/README.md
diff --git a/rms-norm/rms_norm.cu → kernels/rms-norm/rms_norm.cu b/rms-norm/rms_norm.cu → kernels/rms-norm/rms_norm.cu
diff --git a/rms-norm/rms_norm.py → kernels/rms-norm/rms_norm.py b/rms-norm/rms_norm.py → kernels/rms-norm/rms_norm.py
diff --git a/rope/.gitignore → kernels/rope/.gitignore b/rope/.gitignore → kernels/rope/.gitignore
diff --git a/rope/README.md → kernels/rope/README.md b/rope/README.md → kernels/rope/README.md
diff --git a/rope/rope.cu → kernels/rope/rope.cu b/rope/rope.cu → kernels/rope/rope.cu
diff --git a/rope/rope.py → kernels/rope/rope.py b/rope/rope.py → kernels/rope/rope.py
diff --git a/sgemm/.gitignore → kernels/sgemm/.gitignore b/sgemm/.gitignore → kernels/sgemm/.gitignore
diff --git a/sgemm/README.md → kernels/sgemm/README.md b/sgemm/README.md → kernels/sgemm/README.md
diff --git a/sgemm/sgemm.cu → kernels/sgemm/sgemm.cu b/sgemm/sgemm.cu → kernels/sgemm/sgemm.cu
diff --git a/sgemm/sgemm.py → kernels/sgemm/sgemm.py b/sgemm/sgemm.py → kernels/sgemm/sgemm.py
diff --git a/sgemm/sgemm_async.cu → kernels/sgemm/sgemm_async.cu b/sgemm/sgemm_async.cu → kernels/sgemm/sgemm_async.cu
diff --git a/sgemm/sgemm_cublas.cu → kernels/sgemm/sgemm_cublas.cu b/sgemm/sgemm_cublas.cu → kernels/sgemm/sgemm_cublas.cu
diff --git a/sgemm/sgemm_wmma_tf32_stage.cu → kernels/sgemm/sgemm_wmma_tf32_stage.cu b/sgemm/sgemm_wmma_tf32_stage.cu → kernels/sgemm/sgemm_wmma_tf32_stage.cu
diff --git a/sgemv/.gitignore → kernels/sgemv/.gitignore b/sgemv/.gitignore → kernels/sgemv/.gitignore
diff --git a/sgemv/README.md → kernels/sgemv/README.md b/sgemv/README.md → kernels/sgemv/README.md
diff --git a/sgemv/sgemv.cu → kernels/sgemv/sgemv.cu b/sgemv/sgemv.cu → kernels/sgemv/sgemv.cu
diff --git a/sgemv/sgemv.py → kernels/sgemv/sgemv.py b/sgemv/sgemv.py → kernels/sgemv/sgemv.py
diff --git a/sigmoid/.gitignore → kernels/sigmoid/.gitignore b/sigmoid/.gitignore → kernels/sigmoid/.gitignore
diff --git a/sigmoid/README.md → kernels/sigmoid/README.md b/sigmoid/README.md → kernels/sigmoid/README.md
diff --git a/sigmoid/sigmoid.cu → kernels/sigmoid/sigmoid.cu b/sigmoid/sigmoid.cu → kernels/sigmoid/sigmoid.cu
diff --git a/sigmoid/sigmoid.py → kernels/sigmoid/sigmoid.py b/sigmoid/sigmoid.py → kernels/sigmoid/sigmoid.py
diff --git a/softmax/.gitignore → kernels/softmax/.gitignore b/softmax/.gitignore → kernels/softmax/.gitignore
diff --git a/softmax/README.md → kernels/softmax/README.md b/softmax/README.md → kernels/softmax/README.md
diff --git a/softmax/softmax.cu → kernels/softmax/softmax.cu b/softmax/softmax.cu → kernels/softmax/softmax.cu
diff --git a/softmax/softmax.py → kernels/softmax/softmax.py b/softmax/softmax.py → kernels/softmax/softmax.py
diff --git a/swish/.gitignore → kernels/swish/.gitignore b/swish/.gitignore → kernels/swish/.gitignore
diff --git a/swish/README.md → kernels/swish/README.md b/swish/README.md → kernels/swish/README.md
diff --git a/swish/swish.cu → kernels/swish/swish.cu b/swish/swish.cu → kernels/swish/swish.cu
diff --git a/swish/swish.py → kernels/swish/swish.py b/swish/swish.py → kernels/swish/swish.py
diff --git a/openai-triton/.gitignore → kernels/transformer/.gitignore b/openai-triton/.gitignore → kernels/transformer/.gitignore
diff --git a/pytorch/.gitignore → others/pytorch/.gitignore b/pytorch/.gitignore → others/pytorch/.gitignore
diff --git a/pytorch/custom_ops/.gitignore → others/pytorch/custom_ops/.gitignore b/pytorch/custom_ops/.gitignore → others/pytorch/custom_ops/.gitignore
diff --git a/pytorch/distributed/.gitignore → others/pytorch/distributed/.gitignore b/pytorch/distributed/.gitignore → others/pytorch/distributed/.gitignore
diff --git a/pytorch/distributed/test_all_gather.py → ...rs/pytorch/distributed/test_all_gather.py b/pytorch/distributed/test_all_gather.py → ...rs/pytorch/distributed/test_all_gather.py
diff --git a/...ch/distributed/test_all_gather_objects.py → ...ch/distributed/test_all_gather_objects.py b/...ch/distributed/test_all_gather_objects.py → ...ch/distributed/test_all_gather_objects.py
diff --git a/pytorch/distributed/test_all_reduce.py → ...rs/pytorch/distributed/test_all_reduce.py b/pytorch/distributed/test_all_reduce.py → ...rs/pytorch/distributed/test_all_reduce.py
diff --git a/pytorch/distributed/test_all_to_all.py → ...rs/pytorch/distributed/test_all_to_all.py b/pytorch/distributed/test_all_to_all.py → ...rs/pytorch/distributed/test_all_to_all.py
diff --git a/pytorch/distributed/test_broadcast.py → others/pytorch/distributed/test_broadcast.py b/pytorch/distributed/test_broadcast.py → others/pytorch/distributed/test_broadcast.py
diff --git a/pytorch/distributed/test_gather.py → others/pytorch/distributed/test_gather.py b/pytorch/distributed/test_gather.py → others/pytorch/distributed/test_gather.py
diff --git a/pytorch/distributed/test_p2p.py → others/pytorch/distributed/test_p2p.py b/pytorch/distributed/test_p2p.py → others/pytorch/distributed/test_p2p.py
diff --git a/pytorch/distributed/test_reduce.py → others/pytorch/distributed/test_reduce.py b/pytorch/distributed/test_reduce.py → others/pytorch/distributed/test_reduce.py
diff --git a/pytorch/distributed/test_reduce_scatter.py → ...ytorch/distributed/test_reduce_scatter.py b/pytorch/distributed/test_reduce_scatter.py → ...ytorch/distributed/test_reduce_scatter.py
diff --git a/pytorch/distributed/test_scatter.py → others/pytorch/distributed/test_scatter.py b/pytorch/distributed/test_scatter.py → others/pytorch/distributed/test_scatter.py
diff --git a/pytorch/slides/pytorch_2.pdf → others/pytorch/slides/pytorch_2.pdf b/pytorch/slides/pytorch_2.pdf → others/pytorch/slides/pytorch_2.pdf
diff --git a/tensorrt/README.md → others/tensorrt/README.md b/tensorrt/README.md → others/tensorrt/README.md
diff --git a/tensorrt/fmha/.gitignore → others/tensorrt/fmha/.gitignore b/tensorrt/fmha/.gitignore → others/tensorrt/fmha/.gitignore
diff --git a/tensorrt/fmha/README.md → others/tensorrt/fmha/README.md b/tensorrt/fmha/README.md → others/tensorrt/fmha/README.md
diff --git a/tensorrt/fmha/export_fmha.py → others/tensorrt/fmha/export_fmha.py b/tensorrt/fmha/export_fmha.py → others/tensorrt/fmha/export_fmha.py
diff --git a/tensorrt/fmha/fmha_pattern_match_ops.py → ...s/tensorrt/fmha/fmha_pattern_match_ops.py b/tensorrt/fmha/fmha_pattern_match_ops.py → ...s/tensorrt/fmha/fmha_pattern_match_ops.py
diff --git a/tensorrt/plugin/.gitignore → others/tensorrt/plugin/.gitignore b/tensorrt/plugin/.gitignore → others/tensorrt/plugin/.gitignore
diff --git a/tensorrt/plugin/README.md → others/tensorrt/plugin/README.md b/tensorrt/plugin/README.md → others/tensorrt/plugin/README.md
diff --git a/transformer/.gitignore → slides/cuda-slides/.gitignore b/transformer/.gitignore → slides/cuda-slides/.gitignore
diff --git a/cuda-slides/CUDA_C_Programming_Guide_125.pdf → ...a-slides/CUDA_C_Programming_Guide_125.pdf b/cuda-slides/CUDA_C_Programming_Guide_125.pdf → ...a-slides/CUDA_C_Programming_Guide_125.pdf
diff --git a/... API, Enhancements, and NVIDIA Hopper.pdf → ... API, Enhancements, and NVIDIA Hopper.pdf b/... API, Enhancements, and NVIDIA Hopper.pdf → ... API, Enhancements, and NVIDIA Hopper.pdf
diff --git a/...ERENCE WITH SPLITK WORK DECOMPOSITION.pdf → ...ERENCE WITH SPLITK WORK DECOMPOSITION.pdf b/...ERENCE WITH SPLITK WORK DECOMPOSITION.pdf → ...ERENCE WITH SPLITK WORK DECOMPOSITION.pdf
diff --git a/...ation in CUTLASS_1647481957785001Hmln.pdf → ...ation in CUTLASS_1647481957785001Hmln.pdf b/...ation in CUTLASS_1647481957785001Hmln.pdf → ...ation in CUTLASS_1647481957785001Hmln.pdf
diff --git a/...PUs via Persistent and Elastic Blocks.pdf → ...PUs via Persistent and Elastic Blocks.pdf b/...PUs via Persistent and Elastic Blocks.pdf → ...PUs via Persistent and Elastic Blocks.pdf
diff --git a/cuda-slides/CUTLASS/GPU Load Balancing.pdf → ...uda-slides/CUTLASS/GPU Load Balancing.pdf b/cuda-slides/CUTLASS/GPU Load Balancing.pdf → ...uda-slides/CUTLASS/GPU Load Balancing.pdf
diff --git a/...es/CUTLASS/Graphene-CUTE-CUTLASS-2023.pdf → ...es/CUTLASS/Graphene-CUTE-CUTLASS-2023.pdf b/...es/CUTLASS/Graphene-CUTE-CUTLASS-2023.pdf → ...es/CUTLASS/Graphene-CUTE-CUTLASS-2023.pdf
diff --git a/...or cuBLAS-like Performance_ a Worklog.pdf → ...or cuBLAS-like Performance_ a Worklog.pdf b/...or cuBLAS-like Performance_ a Worklog.pdf → ...or cuBLAS-like Performance_ a Worklog.pdf
diff --git a/...ions on Nvidia GPUs using Tensor Core.pdf → ...ions on Nvidia GPUs using Tensor Core.pdf b/...ions on Nvidia GPUs using Tensor Core.pdf → ...ions on Nvidia GPUs using Tensor Core.pdf
diff --git a/...grammability, Performance & Precision.pdf → ...grammability, Performance & Precision.pdf b/...grammability, Performance & Precision.pdf → ...grammability, Performance & Precision.pdf
diff --git a/...Cores_in_CUTLASS_1617596903236001ZNDn.pdf → ...Cores_in_CUTLASS_1617596903236001ZNDn.pdf b/...Cores_in_CUTLASS_1617596903236001ZNDn.pdf → ...Cores_in_CUTLASS_1617596903236001ZNDn.pdf
diff --git a/...es/CUTLASS/colfax-gemm-kernels-hopper.pdf → ...es/CUTLASS/colfax-gemm-kernels-hopper.pdf b/...es/CUTLASS/colfax-gemm-kernels-hopper.pdf → ...es/CUTLASS/colfax-gemm-kernels-hopper.pdf
diff --git a/cuda-slides/CUTLASS/layout_algebra.pdf → ...es/cuda-slides/CUTLASS/layout_algebra.pdf b/cuda-slides/CUTLASS/layout_algebra.pdf → ...es/cuda-slides/CUTLASS/layout_algebra.pdf
diff --git a/...-to-the-absolute-limit-on-nvidia-a100.pdf → ...-to-the-absolute-limit-on-nvidia-a100.pdf b/...-to-the-absolute-limit-on-nvidia-a100.pdf → ...-to-the-absolute-limit-on-nvidia-a100.pdf
diff --git a/...ommunication Techniques and Libraries.pdf → ...ommunication Techniques and Libraries.pdf b/...ommunication Techniques and Libraries.pdf → ...ommunication Techniques and Libraries.pdf
diff --git a/cuda-slides/Graphene-NV-Tensor-IR.pdf → slides/cuda-slides/Graphene-NV-Tensor-IR.pdf b/cuda-slides/Graphene-NV-Tensor-IR.pdf → slides/cuda-slides/Graphene-NV-Tensor-IR.pdf
diff --git a/cuda-slides/Hopper_Tuning_Guide.pdf → slides/cuda-slides/Hopper_Tuning_Guide.pdf b/cuda-slides/Hopper_Tuning_Guide.pdf → slides/cuda-slides/Hopper_Tuning_Guide.pdf
diff --git a/cuda-slides/Inline_PTX_Assembly.pdf → slides/cuda-slides/Inline_PTX_Assembly.pdf b/cuda-slides/Inline_PTX_Assembly.pdf → slides/cuda-slides/Inline_PTX_Assembly.pdf
diff --git a/cuda-slides/MULTI-GPU TRAINING WITHNCCL.pdf → ...da-slides/MULTI-GPU TRAINING WITHNCCL.pdf b/cuda-slides/MULTI-GPU TRAINING WITHNCCL.pdf → ...da-slides/MULTI-GPU TRAINING WITHNCCL.pdf
diff --git a/cuda-slides/NCCL 2.0.pdf → slides/cuda-slides/NCCL 2.0.pdf b/cuda-slides/NCCL 2.0.pdf → slides/cuda-slides/NCCL 2.0.pdf
diff --git a/...Tensor Core GPU Architecture Overview.pdf → ...Tensor Core GPU Architecture Overview.pdf b/...Tensor Core GPU Architecture Overview.pdf → ...Tensor Core GPU Architecture Overview.pdf
diff --git a/cuda-slides/NVIDIA-Torch-TensorRT.pdf → slides/cuda-slides/NVIDIA-Torch-TensorRT.pdf b/cuda-slides/NVIDIA-Torch-TensorRT.pdf → slides/cuda-slides/NVIDIA-Torch-TensorRT.pdf
diff --git a/...ling Argonne National Labs 2022-06-30.pdf → ...ling Argonne National Labs 2022-06-30.pdf b/...ling Argonne National Labs 2022-06-30.pdf → ...ling Argonne National Labs 2022-06-30.pdf
diff --git a/cuda-slides/ORT-Python-Docs.pdf → slides/cuda-slides/ORT-Python-Docs.pdf b/cuda-slides/ORT-Python-Docs.pdf → slides/cuda-slides/ORT-Python-Docs.pdf
diff --git a/...COMMUNICATIONFOR LARGE-SCALE TRAINING.pdf → ...COMMUNICATIONFOR LARGE-SCALE TRAINING.pdf b/...COMMUNICATIONFOR LARGE-SCALE TRAINING.pdf → ...COMMUNICATIONFOR LARGE-SCALE TRAINING.pdf
diff --git a/...L FORDEEP LEARNING TRAINING, AND MORE.pdf → ...L FORDEEP LEARNING TRAINING, AND MORE.pdf b/...L FORDEEP LEARNING TRAINING, AND MORE.pdf → ...L FORDEEP LEARNING TRAINING, AND MORE.pdf
diff --git a/cuda-slides/TORCH.FX.pdf → slides/cuda-slides/TORCH.FX.pdf b/cuda-slides/TORCH.FX.pdf → slides/cuda-slides/TORCH.FX.pdf
diff --git a/cuda-slides/TensorRT-API.pdf → slides/cuda-slides/TensorRT-API.pdf b/cuda-slides/TensorRT-API.pdf → slides/cuda-slides/TensorRT-API.pdf
diff --git a/...-slides/TensorRT-Developer-Guide 10.1.pdf → ...-slides/TensorRT-Developer-Guide 10.1.pdf b/...-slides/TensorRT-Developer-Guide 10.1.pdf → ...-slides/TensorRT-Developer-Guide 10.1.pdf
diff --git a/cuda-slides/TensorRT-Operators.pdf → slides/cuda-slides/TensorRT-Operators.pdf b/cuda-slides/TensorRT-Operators.pdf → slides/cuda-slides/TensorRT-Operators.pdf
diff --git a/cuda-slides/gtc22-whitepaper-hopper.pdf → ...s/cuda-slides/gtc22-whitepaper-hopper.pdf b/cuda-slides/gtc22-whitepaper-hopper.pdf → ...s/cuda-slides/gtc22-whitepaper-hopper.pdf
diff --git a/...nvidia-ampere-architecture-whitepaper.pdf → ...nvidia-ampere-architecture-whitepaper.pdf b/...nvidia-ampere-architecture-whitepaper.pdf → ...nvidia-ampere-architecture-whitepaper.pdf
diff --git a/cuda-slides/ptx_isa_8.5.pdf → slides/cuda-slides/ptx_isa_8.5.pdf b/cuda-slides/ptx_isa_8.5.pdf → slides/cuda-slides/ptx_isa_8.5.pdf
diff --git a/cuda-slides/pytorch-nsys-profiling.pdf → ...es/cuda-slides/pytorch-nsys-profiling.pdf b/cuda-slides/pytorch-nsys-profiling.pdf → ...es/cuda-slides/pytorch-nsys-profiling.pdf
diff --git a/cuda-slides/s21417-faster-transformer.pdf → ...cuda-slides/s21417-faster-transformer.pdf b/cuda-slides/s21417-faster-transformer.pdf → ...cuda-slides/s21417-faster-transformer.pdf
diff --git a/...-slides/volta-architecture-whitepaper.pdf → ...-slides/volta-architecture-whitepaper.pdf b/...-slides/volta-architecture-whitepaper.pdf → ...-slides/volta-architecture-whitepaper.pdf
diff --git a/... Project Update @ Second vLLM Meetup.pptx → ... Project Update @ Second vLLM Meetup.pptx b/... Project Update @ Second vLLM Meetup.pptx → ... Project Update @ Second vLLM Meetup.pptx
diff --git a/vllm-slides/blogs/README.md → slides/vllm-slides/blogs/README.md b/vllm-slides/blogs/README.md → slides/vllm-slides/blogs/README.md
diff --git a/...logs/vllm-automatic-prefix-caching.drawio → ...logs/vllm-automatic-prefix-caching.drawio b/...logs/vllm-automatic-prefix-caching.drawio → ...logs/vllm-automatic-prefix-caching.drawio
diff --git a/.../vllm-automatic-prefix-caching.drawio.png → .../vllm-automatic-prefix-caching.drawio.png b/.../vllm-automatic-prefix-caching.drawio.png → .../vllm-automatic-prefix-caching.drawio.png
diff --git a/...m-prefix-prefill-triton-kernel-tiling.png → ...m-prefix-prefill-triton-kernel-tiling.png b/...m-prefix-prefill-triton-kernel-tiling.png → ...m-prefix-prefill-triton-kernel-tiling.png
diff --git a/.../vllm-prefix-prefill-triton-kernel.drawio → .../vllm-prefix-prefill-triton-kernel.drawio b/.../vllm-prefix-prefill-triton-kernel.drawio → .../vllm-prefix-prefill-triton-kernel.drawio
diff --git a/...vLLM First SF Meetup Slides (Public).pptx → ...vLLM First SF Meetup Slides (Public).pptx b/...vLLM First SF Meetup Slides (Public).pptx → ...vLLM First SF Meetup Slides (Public).pptx
diff --git a/... Update @ Third vLLM Meetup (Public).pptx → ... Update @ Third vLLM Meetup (Public).pptx b/... Update @ Third vLLM Meetup (Public).pptx → ... Update @ Third vLLM Meetup (Public).pptx
diff --git a/third-party/.gitignore b/third-party/.gitignore
@@ -7,4 +7,6 @@
 build
 *.whl
 tmp
+bin
+