From a23b7e85ad0e38c1454c55e27056972574fe80a0 Mon Sep 17 00:00:00 2001 From: wnma3mz Date: Thu, 21 Dec 2023 09:24:30 +0800 Subject: [PATCH] add attention fuse --- 1970/01/01/hello-world/index.html | 6 +- .../index.html" | 6 +- .../index.html" | 10 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../\350\277\220\347\273\264/ssh/index.html" | 10 +- .../index.html" | 10 +- .../index.html" | 14 +- .../index.html" | 14 +- .../index.html" | 10 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../30/\346\235\202\350\256\260/index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 10 +- .../index.html" | 6 +- .../ipv6\345\205\215\346\265\201/index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- 2021/03/18/Python/Birthday_in_Pi/index.html | 6 +- .../index.html" | 6 +- .../Mongodb\344\272\214/index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html | 6 +- .../vLLM\347\232\204PageAttention/index.html" | 10 +- .../index.html" | 10 +- .../index.html" | 14 +- .../index.html" | 14 +- .../PyTorch/pytorch_cpp_extension/index.html | 188 +- .../index.html" | 6 +- .../index.html" | 6 +- .../index.html" | 816 ++ archives/1970/01/index.html | 6 +- archives/1970/index.html | 6 +- archives/2017/04/index.html | 6 +- archives/2017/05/index.html | 6 +- archives/2017/08/index.html | 6 +- archives/2017/11/index.html | 6 +- archives/2017/index.html | 6 +- archives/2018/01/index.html | 6 +- archives/2018/02/index.html | 6 +- archives/2018/03/index.html | 14 +- archives/2018/04/index.html | 6 +- archives/2018/05/index.html | 6 +- archives/2018/08/index.html | 6 +- archives/2018/09/index.html | 6 +- archives/2018/index.html | 8 +- archives/2018/page/2/index.html | 16 +- archives/2018/page/3/index.html | 8 +- archives/2019/05/index.html | 6 +- archives/2019/09/index.html | 6 +- archives/2019/index.html | 6 +- archives/2020/06/index.html | 6 +- archives/2020/08/index.html | 6 +- archives/2020/09/index.html | 6 +- archives/2020/index.html | 6 +- archives/2021/02/index.html | 6 +- archives/2021/03/index.html | 6 +- archives/2021/07/index.html | 6 +- archives/2021/10/index.html | 6 +- archives/2021/index.html | 6 +- archives/2023/01/index.html | 6 +- archives/2023/04/index.html | 6 +- archives/2023/08/index.html | 6 +- archives/2023/09/index.html | 6 +- archives/2023/11/index.html | 14 +- archives/2023/12/index.html | 26 +- archives/2023/index.html | 56 +- archives/2023/page/2/index.html | 28 +- archives/index.html | 56 +- archives/page/2/index.html | 48 +- archives/page/3/index.html | 48 +- archives/page/4/index.html | 51 +- archives/page/5/index.html | 56 +- archives/page/6/index.html | 48 +- archives/page/7/index.html | 34 +- categories/Algorithm/index.html | 4 +- categories/Android/index.html | 4 +- categories/CUDA/index.html | 4 +- categories/Database/index.html | 4 +- categories/DeepLearning/index.html | 24 +- categories/Environment/index.html | 4 +- categories/Life/index.html | 4 +- categories/Linux/index.html | 4 +- categories/MachineLearning/index.html | 4 +- categories/NLP/index.html | 4 +- categories/Note/index.html | 4 +- categories/PaperReading/index.html | 4 +- categories/Python/index.html | 12 +- categories/Raspberry/index.html | 4 +- categories/Scrapy/index.html | 4 +- categories/Software/index.html | 4 +- categories/Visualization/index.html | 4 +- categories/index.html | 6 +- categories/test/index.html | 4 +- css/main.css | 2 +- index.html | 153 +- page/2/index.html | 142 +- page/3/index.html | 140 +- page/4/index.html | 148 +- page/5/index.html | 134 +- page/6/index.html | 142 +- page/7/index.html | 90 +- search.xml | 9327 +++++++++-------- sitemap.xml | 405 +- tags/Algorithm/index.html | 4 +- tags/Android/index.html | 4 +- tags/Attention/index.html | 4 +- tags/CUDA/index.html | 24 +- tags/CUDNN/index.html | 4 +- tags/Cpp/index.html | 24 +- tags/Decorators/index.html | 4 +- tags/DroidCam/index.html | 4 +- tags/GIL/index.html | 4 +- tags/GPU/index.html | 4 +- tags/Github/index.html | 4 +- tags/Google-Charts/index.html | 4 +- tags/HTTPS/index.html | 4 +- tags/Hadoop/index.html | 4 +- tags/HashMap/index.html | 4 +- tags/Hive/index.html | 4 +- tags/IOS/index.html | 4 +- tags/Interpretability/index.html | 4 +- tags/Kindle/index.html | 4 +- tags/Linux/index.html | 6 +- tags/Linux/page/2/index.html | 6 +- tags/NLP/index.html | 4 +- tags/NLTK/index.html | 4 +- tags/Nginx/index.html | 4 +- tags/OpenWrt/index.html | 4 +- tags/Python/page/2/index.html | 6 +- tags/Pytorch/index.html | 24 +- tags/Raspberry/index.html | 4 +- tags/SQL/index.html | 4 +- tags/Scripts/index.html | 4 +- tags/Steaming/index.html | 4 +- tags/Tailscale/index.html | 4 +- tags/Termux/index.html | 4 +- tags/Tmux/index.html | 4 +- tags/Ubuntu/index.html | 4 +- tags/Webhook/index.html | 4 +- tags/adb/index.html | 4 +- tags/caffe/index.html | 4 +- tags/crontab/index.html | 4 +- tags/firewall/index.html | 4 +- tags/ics/index.html | 4 +- tags/importlib/index.html | 4 +- tags/index.html | 6 +- tags/iptables/index.html | 4 +- tags/ipv6/index.html | 4 +- tags/magic-methods/index.html | 4 +- tags/matplotlib/index.html | 4 +- tags/mongodb/index.html | 4 +- tags/netsh/index.html | 4 +- tags/note/index.html | 4 +- tags/numpy/index.html | 4 +- tags/optimization/index.html | 4 +- tags/pandas/index.html | 4 +- tags/python/index.html | 6 +- tags/scheduler/index.html | 4 +- tags/ssh/index.html | 4 +- tags/tensorflow/index.html | 4 +- tags/test/index.html | 4 +- tags/ufw/index.html | 4 +- tags/vLLM/index.html | 4 +- tags/windows/index.html | 4 +- .../index.html" | 4 +- .../index.html" | 4 +- "tags/\345\205\203\347\261\273/index.html" | 4 +- .../index.html" | 4 +- .../index.html" | 4 +- .../index.html" | 4 +- .../index.html" | 4 +- .../index.html" | 4 +- "tags/\345\276\256\344\277\241/index.html" | 4 +- .../index.html" | 4 +- .../index.html" | 4 +- "tags/\346\225\260\345\255\246/index.html" | 4 +- .../index.html" | 4 +- .../index.html" | 4 +- .../index.html" | 4 +- .../index.html" | 4 +- "tags/\347\210\254\350\231\253/index.html" | 4 +- .../index.html" | 4 +- .../index.html" | 4 +- .../index.html" | 4 +- "tags/\347\254\224\350\256\260/index.html" | 4 +- .../index.html" | 4 +- .../index.html" | 4 +- 219 files changed, 7154 insertions(+), 6150 deletions(-) create mode 100644 "2023/12/20/PyTorch/\344\273\245Attention\344\270\272\344\276\213\347\232\204\347\256\227\345\255\220\350\236\215\345\220\210/index.html" diff --git a/1970/01/01/hello-world/index.html b/1970/01/01/hello-world/index.html index a42c4efcf..91d032cea 100644 --- a/1970/01/01/hello-world/index.html +++ b/1970/01/01/hello-world/index.html @@ -90,7 +90,7 @@
-
+
@@ -229,7 +229,7 @@

更新于 - + @@ -380,7 +380,7 @@

Deploy to remote sites

- 61 + 62 日志
diff --git "a/2017/04/09/\347\216\257\345\242\203\345\256\211\350\243\205/Hadoop\344\274\252\345\210\206\345\270\203\345\274\217\345\256\211\350\243\205\357\274\210Centos7\357\274\211/index.html" "b/2017/04/09/\347\216\257\345\242\203\345\256\211\350\243\205/Hadoop\344\274\252\345\210\206\345\270\203\345\274\217\345\256\211\350\243\205\357\274\210Centos7\357\274\211/index.html" index 5f0f348b3..a9370d89c 100644 --- "a/2017/04/09/\347\216\257\345\242\203\345\256\211\350\243\205/Hadoop\344\274\252\345\210\206\345\270\203\345\274\217\345\256\211\350\243\205\357\274\210Centos7\357\274\211/index.html" +++ "b/2017/04/09/\347\216\257\345\242\203\345\256\211\350\243\205/Hadoop\344\274\252\345\210\206\345\270\203\345\274\217\345\256\211\350\243\205\357\274\210Centos7\357\274\211/index.html" @@ -94,7 +94,7 @@
-
+
@@ -233,7 +233,7 @@

更新于 - + @@ -408,7 +408,7 @@

HDFS的常用操作命令

- 61 + 62 日志
diff --git "a/2017/05/07/\347\216\257\345\242\203\345\256\211\350\243\205/\345\234\250Hadoop\345\237\272\347\241\200\344\270\212Hive\347\232\204\345\256\211\350\243\205/index.html" "b/2017/05/07/\347\216\257\345\242\203\345\256\211\350\243\205/\345\234\250Hadoop\345\237\272\347\241\200\344\270\212Hive\347\232\204\345\256\211\350\243\205/index.html" index 9e101151f..a80e3ee4a 100644 --- "a/2017/05/07/\347\216\257\345\242\203\345\256\211\350\243\205/\345\234\250Hadoop\345\237\272\347\241\200\344\270\212Hive\347\232\204\345\256\211\350\243\205/index.html" +++ "b/2017/05/07/\347\216\257\345\242\203\345\256\211\350\243\205/\345\234\250Hadoop\345\237\272\347\241\200\344\270\212Hive\347\232\204\345\256\211\350\243\205/index.html" @@ -32,8 +32,8 @@ - + @@ -94,7 +94,7 @@
-
+
@@ -233,7 +233,7 @@

更新于 - + @@ -369,8 +369,8 @@

转换类型

- +
@@ -460,7 +460,7 @@

转换类型

- 61 + 62 日志
diff --git "a/2017/08/14/\347\216\257\345\242\203\345\256\211\350\243\205/\351\243\216\346\240\274\350\277\201\347\247\273\350\275\254\346\215\242\347\232\204\344\270\244\347\247\215\345\256\236\347\216\260/index.html" "b/2017/08/14/\347\216\257\345\242\203\345\256\211\350\243\205/\351\243\216\346\240\274\350\277\201\347\247\273\350\275\254\346\215\242\347\232\204\344\270\244\347\247\215\345\256\236\347\216\260/index.html" index 5a4f92b2f..a75808a59 100644 --- "a/2017/08/14/\347\216\257\345\242\203\345\256\211\350\243\205/\351\243\216\346\240\274\350\277\201\347\247\273\350\275\254\346\215\242\347\232\204\344\270\244\347\247\215\345\256\236\347\216\260/index.html" +++ "b/2017/08/14/\347\216\257\345\242\203\345\256\211\350\243\205/\351\243\216\346\240\274\350\277\201\347\247\273\350\275\254\346\215\242\347\232\204\344\270\244\347\247\215\345\256\236\347\216\260/index.html" @@ -93,7 +93,7 @@
-
+
@@ -232,7 +232,7 @@

更新于 - + @@ -503,7 +503,7 @@

总结

- 61 + 62 日志
diff --git "a/2017/08/24/\350\277\220\347\273\264/crontab\345\221\275\344\273\244/index.html" "b/2017/08/24/\350\277\220\347\273\264/crontab\345\221\275\344\273\244/index.html" index f342a5467..4b7b7e32e 100644 --- "a/2017/08/24/\350\277\220\347\273\264/crontab\345\221\275\344\273\244/index.html" +++ "b/2017/08/24/\350\277\220\347\273\264/crontab\345\221\275\344\273\244/index.html" @@ -91,7 +91,7 @@
-
+
@@ -230,7 +230,7 @@

更新于 - + @@ -388,7 +388,7 @@

- 61 + 62 日志
diff --git "a/2017/08/26/PythonPackages/Matplotlib\347\254\224\350\256\260/index.html" "b/2017/08/26/PythonPackages/Matplotlib\347\254\224\350\256\260/index.html" index 99b98c5c9..5d6879016 100644 --- "a/2017/08/26/PythonPackages/Matplotlib\347\254\224\350\256\260/index.html" +++ "b/2017/08/26/PythonPackages/Matplotlib\347\254\224\350\256\260/index.html" @@ -91,7 +91,7 @@
-
+
@@ -230,7 +230,7 @@

更新于 - + @@ -391,7 +391,7 @@

搬运工系列

- 61 + 62 日志
diff --git "a/2017/11/18/\350\256\260\344\270\200\346\254\241\345\276\256\344\277\241\345\205\254\344\274\227\345\217\267\347\210\254\350\231\253\347\232\204\347\273\217\345\216\206\357\274\210\345\276\256\344\277\241\346\226\207\347\253\240\351\230\205\350\257\273\347\202\271\350\265\236\347\232\204\350\216\267\345\217\226\357\274\211/index.html" "b/2017/11/18/\350\256\260\344\270\200\346\254\241\345\276\256\344\277\241\345\205\254\344\274\227\345\217\267\347\210\254\350\231\253\347\232\204\347\273\217\345\216\206\357\274\210\345\276\256\344\277\241\346\226\207\347\253\240\351\230\205\350\257\273\347\202\271\350\265\236\347\232\204\350\216\267\345\217\226\357\274\211/index.html" index fb6d3a2bf..f5dab995a 100644 --- "a/2017/11/18/\350\256\260\344\270\200\346\254\241\345\276\256\344\277\241\345\205\254\344\274\227\345\217\267\347\210\254\350\231\253\347\232\204\347\273\217\345\216\206\357\274\210\345\276\256\344\277\241\346\226\207\347\253\240\351\230\205\350\257\273\347\202\271\350\265\236\347\232\204\350\216\267\345\217\226\357\274\211/index.html" +++ "b/2017/11/18/\350\256\260\344\270\200\346\254\241\345\276\256\344\277\241\345\205\254\344\274\227\345\217\267\347\210\254\350\231\253\347\232\204\347\273\217\345\216\206\357\274\210\345\276\256\344\277\241\346\226\207\347\253\240\351\230\205\350\257\273\347\202\271\350\265\236\347\232\204\350\216\267\345\217\226\357\274\211/index.html" @@ -100,7 +100,7 @@
-
+
@@ -239,7 +239,7 @@

更新于 - + @@ -471,7 +471,7 @@

写在最后

- 61 + 62 日志
diff --git "a/2017/11/24/\346\225\260\346\215\256\345\272\223/Mongodb\344\273\216\345\256\211\350\243\205\345\210\260\345\237\272\346\234\254\346\223\215\344\275\234/index.html" "b/2017/11/24/\346\225\260\346\215\256\345\272\223/Mongodb\344\273\216\345\256\211\350\243\205\345\210\260\345\237\272\346\234\254\346\223\215\344\275\234/index.html" index 032aa2983..06abbb384 100644 --- "a/2017/11/24/\346\225\260\346\215\256\345\272\223/Mongodb\344\273\216\345\256\211\350\243\205\345\210\260\345\237\272\346\234\254\346\223\215\344\275\234/index.html" +++ "b/2017/11/24/\346\225\260\346\215\256\345\272\223/Mongodb\344\273\216\345\256\211\350\243\205\345\210\260\345\237\272\346\234\254\346\223\215\344\275\234/index.html" @@ -94,7 +94,7 @@
-
+
@@ -233,7 +233,7 @@

更新于 - + @@ -542,7 +542,7 @@

数据库权限说明表

- 61 + 62 日志
diff --git "a/2018/01/25/IOT/\347\224\250ADB\350\260\203\350\257\225\345\256\211\345\215\223\346\211\213\346\234\272/index.html" "b/2018/01/25/IOT/\347\224\250ADB\350\260\203\350\257\225\345\256\211\345\215\223\346\211\213\346\234\272/index.html" index 093320ddf..2100f53b4 100644 --- "a/2018/01/25/IOT/\347\224\250ADB\350\260\203\350\257\225\345\256\211\345\215\223\346\211\213\346\234\272/index.html" +++ "b/2018/01/25/IOT/\347\224\250ADB\350\260\203\350\257\225\345\256\211\345\215\223\346\211\213\346\234\272/index.html" @@ -95,7 +95,7 @@
-
+
@@ -234,7 +234,7 @@

更新于 - + @@ -428,7 +428,7 @@

使用无线网调试设备

- 61 + 62 日志
diff --git "a/2018/01/30/Python/Python\350\243\205\351\245\260\345\231\250/index.html" "b/2018/01/30/Python/Python\350\243\205\351\245\260\345\231\250/index.html" index d70826d74..32aa8874f 100644 --- "a/2018/01/30/Python/Python\350\243\205\351\245\260\345\231\250/index.html" +++ "b/2018/01/30/Python/Python\350\243\205\351\245\260\345\231\250/index.html" @@ -90,7 +90,7 @@
-
+
@@ -229,7 +229,7 @@

更新于 - + @@ -400,7 +400,7 @@

@property

- 61 + 62 日志
diff --git "a/2018/02/11/\350\277\220\347\273\264/ssh/index.html" "b/2018/02/11/\350\277\220\347\273\264/ssh/index.html" index d0181967c..722a98b8a 100644 --- "a/2018/02/11/\350\277\220\347\273\264/ssh/index.html" +++ "b/2018/02/11/\350\277\220\347\273\264/ssh/index.html" @@ -31,8 +31,8 @@ - + @@ -93,7 +93,7 @@
-
+
@@ -232,7 +232,7 @@

更新于 - + @@ -353,8 +353,8 @@

注意事项

- +
@@ -445,7 +445,7 @@

注意事项

- 61 + 62 日志
diff --git "a/2018/03/02/PythonPackages/Numpy\347\232\20470\344\270\252\346\223\215\344\275\234/index.html" "b/2018/03/02/PythonPackages/Numpy\347\232\20470\344\270\252\346\223\215\344\275\234/index.html" index 9c51845a9..ce10b9867 100644 --- "a/2018/03/02/PythonPackages/Numpy\347\232\20470\344\270\252\346\223\215\344\275\234/index.html" +++ "b/2018/03/02/PythonPackages/Numpy\347\232\20470\344\270\252\346\223\215\344\275\234/index.html" @@ -91,7 +91,7 @@
-
+
@@ -230,7 +230,7 @@

更新于 - + @@ -311,8 +311,8 @@

61-70

ssh笔记（端口转发）

-
@@ -388,7 +388,7 @@

61-70

- 61 + 62 日志
diff --git "a/2018/03/14/Python/Python\346\235\202\350\260\210/index.html" "b/2018/03/14/Python/Python\346\235\202\350\260\210/index.html" index d37147283..db70cbf85 100644 --- "a/2018/03/14/Python/Python\346\235\202\350\260\210/index.html" +++ "b/2018/03/14/Python/Python\346\235\202\350\260\210/index.html" @@ -91,7 +91,7 @@
-
+
@@ -230,7 +230,7 @@

更新于 - + @@ -342,12 +342,12 @@

深刻理解Python中的元类(
-
-

@@ -423,7 +423,7 @@

深刻理解Python中的元类(
- 61 + 62 日志
diff --git "a/2018/03/14/Python/Python\351\255\224\346\234\257\346\226\271\346\263\225/index.html" "b/2018/03/14/Python/Python\351\255\224\346\234\257\346\226\271\346\263\225/index.html" index a16393673..30b4e4317 100644 --- "a/2018/03/14/Python/Python\351\255\224\346\234\257\346\226\271\346\263\225/index.html" +++ "b/2018/03/14/Python/Python\351\255\224\346\234\257\346\226\271\346\263\225/index.html" @@ -90,7 +90,7 @@
-
+
@@ -229,7 +229,7 @@

更新于 - + @@ -365,12 +365,12 @@

__slots__

-
-
@@ -446,7 +446,7 @@

__slots__

- 61 + 62 日志
diff --git "a/2018/03/17/PythonPackages/\344\274\230\345\214\226pandas\345\206\205\345\255\230\345\205\245\351\227\250/index.html" "b/2018/03/17/PythonPackages/\344\274\230\345\214\226pandas\345\206\205\345\255\230\345\205\245\351\227\250/index.html" index 4793e6ffc..a92ec5b86 100644 --- "a/2018/03/17/PythonPackages/\344\274\230\345\214\226pandas\345\206\205\345\255\230\345\205\245\351\227\250/index.html" +++ "b/2018/03/17/PythonPackages/\344\274\230\345\214\226pandas\345\206\205\345\255\230\345\205\245\351\227\250/index.html" @@ -91,7 +91,7 @@
-
+
@@ -230,7 +230,7 @@

更新于 - + @@ -385,8 +385,8 @@

总结

-
- 61 + 62 日志
diff --git "a/2018/04/05/\344\274\240\347\273\237\347\256\227\346\263\225/\346\265\267\351\207\217\346\225\260\346\215\256\345\244\204\347\220\206\351\227\256\351\242\230\357\274\210HashMap+\345\210\206\346\262\273\346\263\225\343\200\201\344\275\215\345\233\276\346\263\225\357\274\211/index.html" "b/2018/04/05/\344\274\240\347\273\237\347\256\227\346\263\225/\346\265\267\351\207\217\346\225\260\346\215\256\345\244\204\347\220\206\351\227\256\351\242\230\357\274\210HashMap+\345\210\206\346\262\273\346\263\225\343\200\201\344\275\215\345\233\276\346\263\225\357\274\211/index.html" index 1005e9f50..be769dc2d 100644 --- "a/2018/04/05/\344\274\240\347\273\237\347\256\227\346\263\225/\346\265\267\351\207\217\346\225\260\346\215\256\345\244\204\347\220\206\351\227\256\351\242\230\357\274\210HashMap+\345\210\206\346\262\273\346\263\225\343\200\201\344\275\215\345\233\276\346\263\225\357\274\211/index.html" +++ "b/2018/04/05/\344\274\240\347\273\237\347\256\227\346\263\225/\346\265\267\351\207\217\346\225\260\346\215\256\345\244\204\347\220\206\351\227\256\351\242\230\357\274\210HashMap+\345\210\206\346\262\273\346\263\225\343\200\201\344\275\215\345\233\276\346\263\225\357\274\211/index.html" @@ -91,7 +91,7 @@
-
+
@@ -230,7 +230,7 @@

更新于 - + @@ -409,7 +409,7 @@

其他

- 61 + 62 日志
diff --git "a/2018/04/07/\346\225\260\346\215\256\345\272\223/SQL\350\257\255\345\217\245\347\232\204\344\274\230\345\214\226/index.html" "b/2018/04/07/\346\225\260\346\215\256\345\272\223/SQL\350\257\255\345\217\245\347\232\204\344\274\230\345\214\226/index.html" index dcd13edb9..eb9772f3b 100644 --- "a/2018/04/07/\346\225\260\346\215\256\345\272\223/SQL\350\257\255\345\217\245\347\232\204\344\274\230\345\214\226/index.html" +++ "b/2018/04/07/\346\225\260\346\215\256\345\272\223/SQL\350\257\255\345\217\245\347\232\204\344\274\230\345\214\226/index.html" @@ -91,7 +91,7 @@
-
+
@@ -230,7 +230,7 @@

更新于 - + @@ -555,7 +555,7 @@

其他优化方法

- 61 + 62 日志
diff --git "a/2018/04/13/\346\225\260\345\255\246/\347\237\251\351\230\265\345\210\206\350\247\243\344\273\213\347\273\215\345\217\212Python\345\256\236\347\216\260\357\274\210QR\343\200\201SVD\343\200\201LU\357\274\211/index.html" "b/2018/04/13/\346\225\260\345\255\246/\347\237\251\351\230\265\345\210\206\350\247\243\344\273\213\347\273\215\345\217\212Python\345\256\236\347\216\260\357\274\210QR\343\200\201SVD\343\200\201LU\357\274\211/index.html" index 88cc91019..d1c4ce29b 100644 --- "a/2018/04/13/\346\225\260\345\255\246/\347\237\251\351\230\265\345\210\206\350\247\243\344\273\213\347\273\215\345\217\212Python\345\256\236\347\216\260\357\274\210QR\343\200\201SVD\343\200\201LU\357\274\211/index.html" +++ "b/2018/04/13/\346\225\260\345\255\246/\347\237\251\351\230\265\345\210\206\350\247\243\344\273\213\347\273\215\345\217\212Python\345\256\236\347\216\260\357\274\210QR\343\200\201SVD\343\200\201LU\357\274\211/index.html" @@ -94,7 +94,7 @@
-
+
@@ -233,7 +233,7 @@

更新于 - + @@ -545,7 +545,7 @@

三种实现方式

- 61 + 62 日志
diff --git "a/2018/04/14/\351\230\205\350\257\273\347\254\224\350\256\260/\347\245\236\347\273\217\347\275\221\347\273\234\347\254\224\350\256\260\342\205\240/index.html" "b/2018/04/14/\351\230\205\350\257\273\347\254\224\350\256\260/\347\245\236\347\273\217\347\275\221\347\273\234\347\254\224\350\256\260\342\205\240/index.html" index 51ef86881..188d8215b 100644 --- "a/2018/04/14/\351\230\205\350\257\273\347\254\224\350\256\260/\347\245\236\347\273\217\347\275\221\347\273\234\347\254\224\350\256\260\342\205\240/index.html" +++ "b/2018/04/14/\351\230\205\350\257\273\347\254\224\350\256\260/\347\245\236\347\273\217\347\275\221\347\273\234\347\254\224\350\256\260\342\205\240/index.html" @@ -93,7 +93,7 @@
-
+
@@ -232,7 +232,7 @@

更新于 - + @@ -473,7 +473,7 @@

改进神经网络v3.0——dropout

- 61 + 62 日志
diff --git "a/2018/04/30/\346\225\260\345\255\246/\346\213\211\346\240\274\346\234\227\346\227\245\344\271\230\345\255\220\346\263\225\345\255\246\344\271\240\347\254\224\350\256\260/index.html" "b/2018/04/30/\346\225\260\345\255\246/\346\213\211\346\240\274\346\234\227\346\227\245\344\271\230\345\255\220\346\263\225\345\255\246\344\271\240\347\254\224\350\256\260/index.html" index 9b5fc4a74..7d5965b62 100644 --- "a/2018/04/30/\346\225\260\345\255\246/\346\213\211\346\240\274\346\234\227\346\227\245\344\271\230\345\255\220\346\263\225\345\255\246\344\271\240\347\254\224\350\256\260/index.html" +++ "b/2018/04/30/\346\225\260\345\255\246/\346\213\211\346\240\274\346\234\227\346\227\245\344\271\230\345\255\220\346\263\225\345\255\246\344\271\240\347\254\224\350\256\260/index.html" @@ -92,7 +92,7 @@
-
+
@@ -231,7 +231,7 @@

更新于 - + @@ -482,7 +482,7 @@

对偶概念的简单介绍

- 61 + 62 日志
diff --git "a/2018/05/05/CUDA/\350\256\260\350\242\253CUDA\346\212\230\350\205\276\346\255\273\345\216\273\346\264\273\346\235\245\347\232\204\351\202\243\345\215\201\345\244\251/index.html" "b/2018/05/05/CUDA/\350\256\260\350\242\253CUDA\346\212\230\350\205\276\346\255\273\345\216\273\346\264\273\346\235\245\347\232\204\351\202\243\345\215\201\345\244\251/index.html" index 8bf880d4a..e837497cc 100644 --- "a/2018/05/05/CUDA/\350\256\260\350\242\253CUDA\346\212\230\350\205\276\346\255\273\345\216\273\346\264\273\346\235\245\347\232\204\351\202\243\345\215\201\345\244\251/index.html" +++ "b/2018/05/05/CUDA/\350\256\260\350\242\253CUDA\346\212\230\350\205\276\346\255\273\345\216\273\346\264\273\346\235\245\347\232\204\351\202\243\345\215\201\345\244\251/index.html" @@ -93,7 +93,7 @@
-
+
@@ -232,7 +232,7 @@

更新于 - + @@ -448,7 +448,7 @@

参考文章

- 61 + 62 日志
diff --git "a/2018/05/13/CUDA/CUDA\347\274\226\347\250\213\345\205\245\351\227\250/index.html" "b/2018/05/13/CUDA/CUDA\347\274\226\347\250\213\345\205\245\351\227\250/index.html" index da64ebdd6..5240d668e 100644 --- "a/2018/05/13/CUDA/CUDA\347\274\226\347\250\213\345\205\245\351\227\250/index.html" +++ "b/2018/05/13/CUDA/CUDA\347\274\226\347\250\213\345\205\245\351\227\250/index.html" @@ -92,7 +92,7 @@
-
+
@@ -231,7 +231,7 @@

更新于 - + @@ -398,7 +398,7 @@

性能剖析

- 61 + 62 日志
diff --git "a/2018/05/13/\351\230\205\350\257\273\347\254\224\350\256\260/NLTK\351\230\205\350\257\273\347\254\224\350\256\260\342\205\240/index.html" "b/2018/05/13/\351\230\205\350\257\273\347\254\224\350\256\260/NLTK\351\230\205\350\257\273\347\254\224\350\256\260\342\205\240/index.html" index ad1eeaa16..6efc97ada 100644 --- "a/2018/05/13/\351\230\205\350\257\273\347\254\224\350\256\260/NLTK\351\230\205\350\257\273\347\254\224\350\256\260\342\205\240/index.html" +++ "b/2018/05/13/\351\230\205\350\257\273\347\254\224\350\256\260/NLTK\351\230\205\350\257\273\347\254\224\350\256\260\342\205\240/index.html" @@ -104,7 +104,7 @@
-
+
@@ -243,7 +243,7 @@

更新于 - + @@ -514,7 +514,7 @@

第二章：获得文本语
- 61 + 62 日志
diff --git "a/2018/05/18/CUDA/CUDA\345\256\230\346\226\271\346\226\207\346\241\243\347\254\224\350\256\260/index.html" "b/2018/05/18/CUDA/CUDA\345\256\230\346\226\271\346\226\207\346\241\243\347\254\224\350\256\260/index.html" index c819af4dc..d98ceb853 100644 --- "a/2018/05/18/CUDA/CUDA\345\256\230\346\226\271\346\226\207\346\241\243\347\254\224\350\256\260/index.html" +++ "b/2018/05/18/CUDA/CUDA\345\256\230\346\226\271\346\226\207\346\241\243\347\254\224\350\256\260/index.html" @@ -90,7 +90,7 @@
-
+
@@ -229,7 +229,7 @@

更新于 - + @@ -386,7 +386,7 @@

- 61 + 62 日志
diff --git "a/2018/05/18/CUDA/CUDA\347\274\226\347\250\213\344\272\214\346\254\241\345\205\245\351\227\250/index.html" "b/2018/05/18/CUDA/CUDA\347\274\226\347\250\213\344\272\214\346\254\241\345\205\245\351\227\250/index.html" index a6624ba1d..a74e729e4 100644 --- "a/2018/05/18/CUDA/CUDA\347\274\226\347\250\213\344\272\214\346\254\241\345\205\245\351\227\250/index.html" +++ "b/2018/05/18/CUDA/CUDA\347\274\226\347\250\213\344\272\214\346\254\241\345\205\245\351\227\250/index.html" @@ -90,7 +90,7 @@
-
+
@@ -229,7 +229,7 @@

更新于 - + @@ -394,7 +394,7 @@

补充说明

- 61 + 62 日志
diff --git "a/2018/05/23/\351\230\205\350\257\273\347\254\224\350\256\260/NLTK\351\230\205\350\257\273\347\254\224\350\256\260\342\205\241/index.html" "b/2018/05/23/\351\230\205\350\257\273\347\254\224\350\256\260/NLTK\351\230\205\350\257\273\347\254\224\350\256\260\342\205\241/index.html" index e9b27112d..6a3d5f0d8 100644 --- "a/2018/05/23/\351\230\205\350\257\273\347\254\224\350\256\260/NLTK\351\230\205\350\257\273\347\254\224\350\256\260\342\205\241/index.html" +++ "b/2018/05/23/\351\230\205\350\257\273\347\254\224\350\256\260/NLTK\351\230\205\350\257\273\347\254\224\350\256\260\342\205\241/index.html" @@ -96,7 +96,7 @@
-
+
@@ -235,7 +235,7 @@

更新于 - + @@ -589,7 +589,7 @@

N-gram标注

- 61 + 62 日志
diff --git "a/2018/05/30/\346\235\202\350\256\260/index.html" "b/2018/05/30/\346\235\202\350\256\260/index.html" index b12e42f9e..ae44e4629 100644 --- "a/2018/05/30/\346\235\202\350\256\260/index.html" +++ "b/2018/05/30/\346\235\202\350\256\260/index.html" @@ -89,7 +89,7 @@
-
+
@@ -228,7 +228,7 @@

更新于 - + @@ -474,7 +474,7 @@

互现信息

- 61 + 62 日志
diff --git "a/2018/08/05/\351\230\205\350\257\273\347\254\224\350\256\260/NLTK\351\230\205\350\257\273\347\254\224\350\256\260\342\205\242/index.html" "b/2018/08/05/\351\230\205\350\257\273\347\254\224\350\256\260/NLTK\351\230\205\350\257\273\347\254\224\350\256\260\342\205\242/index.html" index bad424b52..6cba50f54 100644 --- "a/2018/08/05/\351\230\205\350\257\273\347\254\224\350\256\260/NLTK\351\230\205\350\257\273\347\254\224\350\256\260\342\205\242/index.html" +++ "b/2018/08/05/\351\230\205\350\257\273\347\254\224\350\256\260/NLTK\351\230\205\350\257\273\347\254\224\350\256\260\342\205\242/index.html" @@ -92,7 +92,7 @@
-
+
@@ -231,7 +231,7 @@

更新于 - + @@ -432,7 +432,7 @@

涉及到的定义

- 61 + 62 日志
diff --git "a/2018/08/18/\344\272\247\345\223\201\350\257\204\350\256\272\346\225\260\346\215\256\346\203\205\346\204\237\345\210\206\346\236\220/index.html" "b/2018/08/18/\344\272\247\345\223\201\350\257\204\350\256\272\346\225\260\346\215\256\346\203\205\346\204\237\345\210\206\346\236\220/index.html" index baa0433dd..4b27652f7 100644 --- "a/2018/08/18/\344\272\247\345\223\201\350\257\204\350\256\272\346\225\260\346\215\256\346\203\205\346\204\237\345\210\206\346\236\220/index.html" +++ "b/2018/08/18/\344\272\247\345\223\201\350\257\204\350\256\272\346\225\260\346\215\256\346\203\205\346\204\237\345\210\206\346\236\220/index.html" @@ -90,7 +90,7 @@
-
+
@@ -229,7 +229,7 @@

更新于 - + @@ -420,7 +420,7 @@

模糊综合评价

- 61 + 62 日志
diff --git "a/2018/08/20/\351\230\205\350\257\273\347\254\224\350\256\260/\347\245\236\347\273\217\347\275\221\347\273\234\347\254\224\350\256\260\342\205\241/index.html" "b/2018/08/20/\351\230\205\350\257\273\347\254\224\350\256\260/\347\245\236\347\273\217\347\275\221\347\273\234\347\254\224\350\256\260\342\205\241/index.html" index aadbce2cb..051ab3e35 100644 --- "a/2018/08/20/\351\230\205\350\257\273\347\254\224\350\256\260/\347\245\236\347\273\217\347\275\221\347\273\234\347\254\224\350\256\260\342\205\241/index.html" +++ "b/2018/08/20/\351\230\205\350\257\273\347\254\224\350\256\260/\347\245\236\347\273\217\347\275\221\347\273\234\347\254\224\350\256\260\342\205\241/index.html" @@ -92,7 +92,7 @@
-
+
@@ -231,7 +231,7 @@

更新于 - + @@ -441,7 +441,7 @@

LSTM（时间递归神经网络）

- 61 + 62 日志
diff --git "a/2018/09/30/\346\234\254\347\247\221\347\247\221\347\240\224\347\232\204\351\202\243\344\272\233\344\272\213/index.html" "b/2018/09/30/\346\234\254\347\247\221\347\247\221\347\240\224\347\232\204\351\202\243\344\272\233\344\272\213/index.html" index e65c5d6b3..a97ee8ac1 100644 --- "a/2018/09/30/\346\234\254\347\247\221\347\247\221\347\240\224\347\232\204\351\202\243\344\272\233\344\272\213/index.html" +++ "b/2018/09/30/\346\234\254\347\247\221\347\247\221\347\240\224\347\232\204\351\202\243\344\272\233\344\272\213/index.html" @@ -89,7 +89,7 @@
-
+
@@ -228,7 +228,7 @@

更新于 - + @@ -375,7 +375,7 @@

本科科研的那些事

- 61 + 62 日志
diff --git "a/2019/05/13/\346\225\260\345\255\246/\347\272\277\344\273\243\347\237\251\351\230\265/index.html" "b/2019/05/13/\346\225\260\345\255\246/\347\272\277\344\273\243\347\237\251\351\230\265/index.html" index a846f2f4b..f68dfde39 100644 --- "a/2019/05/13/\346\225\260\345\255\246/\347\272\277\344\273\243\347\237\251\351\230\265/index.html" +++ "b/2019/05/13/\346\225\260\345\255\246/\347\272\277\344\273\243\347\237\251\351\230\265/index.html" @@ -90,7 +90,7 @@
-
+
@@ -229,7 +229,7 @@

更新于 - + @@ -401,7 +401,7 @@

- 61 + 62 日志
diff --git "a/2019/05/28/\350\256\272\346\226\207\344\273\216\351\233\266\345\210\260\344\270\200/index.html" "b/2019/05/28/\350\256\272\346\226\207\344\273\216\351\233\266\345\210\260\344\270\200/index.html" index fdb10d3de..8be210f73 100644 --- "a/2019/05/28/\350\256\272\346\226\207\344\273\216\351\233\266\345\210\260\344\270\200/index.html" +++ "b/2019/05/28/\350\256\272\346\226\207\344\273\216\351\233\266\345\210\260\344\270\200/index.html" @@ -89,7 +89,7 @@
-
+
@@ -228,7 +228,7 @@

更新于 - + @@ -437,7 +437,7 @@

Mathpix Snip

- 61 + 62 日志
diff --git "a/2019/09/14/IOT/\345\257\271\345\220\204\347\247\215\347\224\265\345\255\220\350\256\276\345\244\207\347\232\204\346\213\223\345\261\225\347\216\251\346\263\225\344\270\216\344\272\222\350\201\224/index.html" "b/2019/09/14/IOT/\345\257\271\345\220\204\347\247\215\347\224\265\345\255\220\350\256\276\345\244\207\347\232\204\346\213\223\345\261\225\347\216\251\346\263\225\344\270\216\344\272\222\350\201\224/index.html" index 319742df6..fb38f6d22 100644 --- "a/2019/09/14/IOT/\345\257\271\345\220\204\347\247\215\347\224\265\345\255\220\350\256\276\345\244\207\347\232\204\346\213\223\345\261\225\347\216\251\346\263\225\344\270\216\344\272\222\350\201\224/index.html" +++ "b/2019/09/14/IOT/\345\257\271\345\220\204\347\247\215\347\224\265\345\255\220\350\256\276\345\244\207\347\232\204\346\213\223\345\261\225\347\216\251\346\263\225\344\270\216\344\272\222\350\201\224/index.html" @@ -95,7 +95,7 @@
-
+
@@ -234,7 +234,7 @@

更新于 - + @@ -427,7 +427,7 @@

树莓派

- 61 + 62 日志
diff --git "a/2020/06/14/Python/Python3\347\232\204Tools-scripts\346\225\264\347\220\206/index.html" "b/2020/06/14/Python/Python3\347\232\204Tools-scripts\346\225\264\347\220\206/index.html" index a87e972b2..4e4fa2676 100644 --- "a/2020/06/14/Python/Python3\347\232\204Tools-scripts\346\225\264\347\220\206/index.html" +++ "b/2020/06/14/Python/Python3\347\232\204Tools-scripts\346\225\264\347\220\206/index.html" @@ -90,7 +90,7 @@
-
+
@@ -229,7 +229,7 @@

更新于 - + @@ -438,7 +438,7 @@

texi2html.py

- 61 + 62 日志
diff --git "a/2020/08/08/\350\277\220\347\273\264/firewall\347\253\257\345\217\243\350\275\254\345\217\221/index.html" "b/2020/08/08/\350\277\220\347\273\264/firewall\347\253\257\345\217\243\350\275\254\345\217\221/index.html" index 619487e60..099f2cbe7 100644 --- "a/2020/08/08/\350\277\220\347\273\264/firewall\347\253\257\345\217\243\350\275\254\345\217\221/index.html" +++ "b/2020/08/08/\350\277\220\347\273\264/firewall\347\253\257\345\217\243\350\275\254\345\217\221/index.html" @@ -31,8 +31,8 @@ - + @@ -96,7 +96,7 @@
-
+
@@ -235,7 +235,7 @@

更新于 - + @@ -321,8 +321,8 @@

实验

- + @@ -416,7 +416,7 @@

实验

- 61 + 62 日志
diff --git "a/2020/09/16/ICS\345\234\250\347\272\277\350\257\276\350\241\250\345\210\266\344\275\234/index.html" "b/2020/09/16/ICS\345\234\250\347\272\277\350\257\276\350\241\250\345\210\266\344\275\234/index.html" index 60ad8098a..a9e9c7ac2 100644 --- "a/2020/09/16/ICS\345\234\250\347\272\277\350\257\276\350\241\250\345\210\266\344\275\234/index.html" +++ "b/2020/09/16/ICS\345\234\250\347\272\277\350\257\276\350\241\250\345\210\266\344\275\234/index.html" @@ -91,7 +91,7 @@
-
+
@@ -230,7 +230,7 @@

更新于 - + @@ -415,7 +415,7 @@

写到最后

- 61 + 62 日志
diff --git "a/2020/09/16/ipv6\345\205\215\346\265\201/index.html" "b/2020/09/16/ipv6\345\205\215\346\265\201/index.html" index 90a341560..b3eee7c6f 100644 --- "a/2020/09/16/ipv6\345\205\215\346\265\201/index.html" +++ "b/2020/09/16/ipv6\345\205\215\346\265\201/index.html" @@ -91,7 +91,7 @@
-
+
@@ -230,7 +230,7 @@

更新于 - + @@ -374,7 +374,7 @@

- 61 + 62 日志
diff --git "a/2021/02/03/PyTorch/torch-scheduler\350\256\260\345\275\225/index.html" "b/2021/02/03/PyTorch/torch-scheduler\350\256\260\345\275\225/index.html" index 09f827e9d..cf0bd8035 100644 --- "a/2021/02/03/PyTorch/torch-scheduler\350\256\260\345\275\225/index.html" +++ "b/2021/02/03/PyTorch/torch-scheduler\350\256\260\345\275\225/index.html" @@ -91,7 +91,7 @@
-
+
@@ -230,7 +230,7 @@

更新于 - + @@ -378,7 +378,7 @@

- 61 + 62 日志
diff --git "a/2021/02/03/\346\225\260\345\255\246/\346\234\200\344\274\230\345\214\226\345\255\246\344\271\240\347\254\224\350\256\260/index.html" "b/2021/02/03/\346\225\260\345\255\246/\346\234\200\344\274\230\345\214\226\345\255\246\344\271\240\347\254\224\350\256\260/index.html" index 01e7c3e70..8898a3b72 100644 --- "a/2021/02/03/\346\225\260\345\255\246/\346\234\200\344\274\230\345\214\226\345\255\246\344\271\240\347\254\224\350\256\260/index.html" +++ "b/2021/02/03/\346\225\260\345\255\246/\346\234\200\344\274\230\345\214\226\345\255\246\344\271\240\347\254\224\350\256\260/index.html" @@ -93,7 +93,7 @@
-
+
@@ -232,7 +232,7 @@

更新于 - + @@ -532,7 +532,7 @@

相关书籍

- 61 + 62 日志
diff --git a/2021/03/18/Python/Birthday_in_Pi/index.html b/2021/03/18/Python/Birthday_in_Pi/index.html index 23a4a4e2c..10b9a1e75 100644 --- a/2021/03/18/Python/Birthday_in_Pi/index.html +++ b/2021/03/18/Python/Birthday_in_Pi/index.html @@ -92,7 +92,7 @@
-
+
@@ -231,7 +231,7 @@

更新于 - + @@ -531,7 +531,7 @@

参考文章

- 61 + 62 日志
diff --git "a/2021/07/21/IOT/\345\256\211\345\215\223\346\211\213\346\234\272\345\210\251\347\224\250Termux\345\256\211\350\243\205Linux/index.html" "b/2021/07/21/IOT/\345\256\211\345\215\223\346\211\213\346\234\272\345\210\251\347\224\250Termux\345\256\211\350\243\205Linux/index.html" index f77187c8f..5b40f2324 100644 --- "a/2021/07/21/IOT/\345\256\211\345\215\223\346\211\213\346\234\272\345\210\251\347\224\250Termux\345\256\211\350\243\205Linux/index.html" +++ "b/2021/07/21/IOT/\345\256\211\345\215\223\346\211\213\346\234\272\345\210\251\347\224\250Termux\345\256\211\350\243\205Linux/index.html" @@ -92,7 +92,7 @@
-
+
@@ -231,7 +231,7 @@

更新于 - + @@ -389,7 +389,7 @@

其他

- 61 + 62 日志
diff --git "a/2021/07/21/\346\225\260\346\215\256\345\272\223/Mongodb\344\272\214/index.html" "b/2021/07/21/\346\225\260\346\215\256\345\272\223/Mongodb\344\272\214/index.html" index 4ffee77a7..550905fd7 100644 --- "a/2021/07/21/\346\225\260\346\215\256\345\272\223/Mongodb\344\272\214/index.html" +++ "b/2021/07/21/\346\225\260\346\215\256\345\272\223/Mongodb\344\272\214/index.html" @@ -91,7 +91,7 @@
-
+
@@ -230,7 +230,7 @@

更新于 - + @@ -397,7 +397,7 @@

主节点选举配置

- 61 + 62 日志
diff --git "a/2021/07/21/\350\277\220\347\273\264/\345\256\236\351\252\214\345\256\244\347\275\221\347\253\231\346\234\215\345\212\241\350\277\201\347\247\273(\351\235\231\346\200\201\347\275\221\351\241\265)/index.html" "b/2021/07/21/\350\277\220\347\273\264/\345\256\236\351\252\214\345\256\244\347\275\221\347\253\231\346\234\215\345\212\241\350\277\201\347\247\273(\351\235\231\346\200\201\347\275\221\351\241\265)/index.html" index ce5236f9e..9de3f581e 100644 --- "a/2021/07/21/\350\277\220\347\273\264/\345\256\236\351\252\214\345\256\244\347\275\221\347\253\231\346\234\215\345\212\241\350\277\201\347\247\273(\351\235\231\346\200\201\347\275\221\351\241\265)/index.html" +++ "b/2021/07/21/\350\277\220\347\273\264/\345\256\236\351\252\214\345\256\244\347\275\221\347\253\231\346\234\215\345\212\241\350\277\201\347\247\273(\351\235\231\346\200\201\347\275\221\351\241\265)/index.html" @@ -94,7 +94,7 @@
-
+
@@ -233,7 +233,7 @@

更新于 - + @@ -397,7 +397,7 @@

Github Hook配置

- 61 + 62 日志
diff --git "a/2021/10/30/\347\216\257\345\242\203\345\256\211\350\243\205/Ubuntu20.04\345\256\211\350\243\205\346\267\261\345\272\246\345\255\246\344\271\240\347\216\257\345\242\203\345\217\212\351\205\215\347\275\256\344\270\252\344\272\272\347\216\257\345\242\203\345\260\217\350\256\260/index.html" "b/2021/10/30/\347\216\257\345\242\203\345\256\211\350\243\205/Ubuntu20.04\345\256\211\350\243\205\346\267\261\345\272\246\345\255\246\344\271\240\347\216\257\345\242\203\345\217\212\351\205\215\347\275\256\344\270\252\344\272\272\347\216\257\345\242\203\345\260\217\350\256\260/index.html" index bbe5e0023..31b7638ec 100644 --- "a/2021/10/30/\347\216\257\345\242\203\345\256\211\350\243\205/Ubuntu20.04\345\256\211\350\243\205\346\267\261\345\272\246\345\255\246\344\271\240\347\216\257\345\242\203\345\217\212\351\205\215\347\275\256\344\270\252\344\272\272\347\216\257\345\242\203\345\260\217\350\256\260/index.html" +++ "b/2021/10/30/\347\216\257\345\242\203\345\256\211\350\243\205/Ubuntu20.04\345\256\211\350\243\205\346\267\261\345\272\246\345\255\246\344\271\240\347\216\257\345\242\203\345\217\212\351\205\215\347\275\256\344\270\252\344\272\272\347\216\257\345\242\203\345\260\217\350\256\260/index.html" @@ -94,7 +94,7 @@
-
+
@@ -233,7 +233,7 @@

更新于 - + @@ -424,7 +424,7 @@

远程连接

- 61 + 62 日志
diff --git "a/2023/01/10/IOT/OpenWrt\350\267\257\347\224\261\345\231\250\350\256\260\345\275\225/index.html" "b/2023/01/10/IOT/OpenWrt\350\267\257\347\224\261\345\231\250\350\256\260\345\275\225/index.html" index 5b38a6245..f197417ef 100644 --- "a/2023/01/10/IOT/OpenWrt\350\267\257\347\224\261\345\231\250\350\256\260\345\275\225/index.html" +++ "b/2023/01/10/IOT/OpenWrt\350\267\257\347\224\261\345\231\250\350\256\260\345\275\225/index.html" @@ -102,7 +102,7 @@
-
+
@@ -241,7 +241,7 @@

更新于 - + @@ -476,7 +476,7 @@

接口

- 61 + 62 日志
diff --git "a/2023/04/03/\345\271\264\345\272\246\346\216\250\350\215\220\350\275\257\344\273\266Tailscale/index.html" "b/2023/04/03/\345\271\264\345\272\246\346\216\250\350\215\220\350\275\257\344\273\266Tailscale/index.html" index 8ce8426e4..b252f20a8 100644 --- "a/2023/04/03/\345\271\264\345\272\246\346\216\250\350\215\220\350\275\257\344\273\266Tailscale/index.html" +++ "b/2023/04/03/\345\271\264\345\272\246\346\216\250\350\215\220\350\275\257\344\273\266Tailscale/index.html" @@ -100,7 +100,7 @@
-
+
@@ -239,7 +239,7 @@

更新于 - + @@ -470,7 +470,7 @@

Step 3

- 61 + 62 日志
diff --git "a/2023/08/13/Tmux\347\232\204\344\275\277\347\224\250\345\222\214\351\205\215\347\275\256/index.html" "b/2023/08/13/Tmux\347\232\204\344\275\277\347\224\250\345\222\214\351\205\215\347\275\256/index.html" index 3539854f1..4edec38ce 100644 --- "a/2023/08/13/Tmux\347\232\204\344\275\277\347\224\250\345\222\214\351\205\215\347\275\256/index.html" +++ "b/2023/08/13/Tmux\347\232\204\344\275\277\347\224\250\345\222\214\351\205\215\347\275\256/index.html" @@ -91,7 +91,7 @@
-
+
@@ -230,7 +230,7 @@

更新于 - + @@ -414,7 +414,7 @@

其他常用配置

- 61 + 62 日志
diff --git "a/2023/08/13/\351\230\205\350\257\273\347\254\224\350\256\260/\346\250\241\344\273\277\346\270\270\346\210\217\342\200\224\342\200\224\346\234\272\345\231\250\344\274\232\346\200\235\350\200\203\345\220\227/index.html" "b/2023/08/13/\351\230\205\350\257\273\347\254\224\350\256\260/\346\250\241\344\273\277\346\270\270\346\210\217\342\200\224\342\200\224\346\234\272\345\231\250\344\274\232\346\200\235\350\200\203\345\220\227/index.html" index 5ffafbbb7..6e8aac3c8 100644 --- "a/2023/08/13/\351\230\205\350\257\273\347\254\224\350\256\260/\346\250\241\344\273\277\346\270\270\346\210\217\342\200\224\342\200\224\346\234\272\345\231\250\344\274\232\346\200\235\350\200\203\345\220\227/index.html" +++ "b/2023/08/13/\351\230\205\350\257\273\347\254\224\350\256\260/\346\250\241\344\273\277\346\270\270\346\210\217\342\200\224\342\200\224\346\234\272\345\231\250\344\274\232\346\200\235\350\200\203\345\220\227/index.html" @@ -89,7 +89,7 @@
-
+
@@ -228,7 +228,7 @@

更新于 - + @@ -427,7 +427,7 @@

参考文献

- 61 + 62 日志
diff --git "a/2023/08/19/SPSA\351\232\217\346\234\272\344\274\230\345\214\226\347\256\227\346\263\225/index.html" "b/2023/08/19/SPSA\351\232\217\346\234\272\344\274\230\345\214\226\347\256\227\346\263\225/index.html" index d970dc66a..c3ae5bba6 100644 --- "a/2023/08/19/SPSA\351\232\217\346\234\272\344\274\230\345\214\226\347\256\227\346\263\225/index.html" +++ "b/2023/08/19/SPSA\351\232\217\346\234\272\344\274\230\345\214\226\347\256\227\346\263\225/index.html" @@ -90,7 +90,7 @@
-
+
@@ -229,7 +229,7 @@

更新于 - + @@ -437,7 +437,7 @@

Q & A

- 61 + 62 日志
diff --git "a/2023/08/28/\344\274\240\347\273\237\347\256\227\346\263\225/\345\255\227\347\254\246\344\270\262/index.html" "b/2023/08/28/\344\274\240\347\273\237\347\256\227\346\263\225/\345\255\227\347\254\246\344\270\262/index.html" index 9ebfb5992..7a3152678 100644 --- "a/2023/08/28/\344\274\240\347\273\237\347\256\227\346\263\225/\345\255\227\347\254\246\344\270\262/index.html" +++ "b/2023/08/28/\344\274\240\347\273\237\347\256\227\346\263\225/\345\255\227\347\254\246\344\270\262/index.html" @@ -90,7 +90,7 @@
-
+
@@ -229,7 +229,7 @@

更新于 - + @@ -404,7 +404,7 @@

模拟/深搜

- 61 + 62 日志
diff --git "a/2023/09/03/PyTorch/\344\272\244\345\217\211\347\206\265\347\232\204\345\217\215\345\220\221\344\274\240\346\222\255 Step by Step/index.html" "b/2023/09/03/PyTorch/\344\272\244\345\217\211\347\206\265\347\232\204\345\217\215\345\220\221\344\274\240\346\222\255 Step by Step/index.html" index 423316b36..9f5c9e877 100644 --- "a/2023/09/03/PyTorch/\344\272\244\345\217\211\347\206\265\347\232\204\345\217\215\345\220\221\344\274\240\346\222\255 Step by Step/index.html" +++ "b/2023/09/03/PyTorch/\344\272\244\345\217\211\347\206\265\347\232\204\345\217\215\345\220\221\344\274\240\346\222\255 Step by Step/index.html" @@ -92,7 +92,7 @@
-
+
@@ -231,7 +231,7 @@

更新于 - + @@ -478,7 +478,7 @@

梯度计算

- 61 + 62 日志
diff --git "a/2023/11/01/Python/\345\244\232\350\277\233\347\250\213\345\206\205\345\255\230\346\263\204\346\274\217\351\227\256\351\242\230\350\256\260\345\275\225/index.html" "b/2023/11/01/Python/\345\244\232\350\277\233\347\250\213\345\206\205\345\255\230\346\263\204\346\274\217\351\227\256\351\242\230\350\256\260\345\275\225/index.html" index 0f16caa33..02a108124 100644 --- "a/2023/11/01/Python/\345\244\232\350\277\233\347\250\213\345\206\205\345\255\230\346\263\204\346\274\217\351\227\256\351\242\230\350\256\260\345\275\225/index.html" +++ "b/2023/11/01/Python/\345\244\232\350\277\233\347\250\213\345\206\205\345\255\230\346\263\204\346\274\217\351\227\256\351\242\230\350\256\260\345\275\225/index.html" @@ -91,7 +91,7 @@
-
+
@@ -230,7 +230,7 @@

更新于 - + @@ -401,7 +401,7 @@

Q & A

- 61 + 62 日志
diff --git a/2023/11/17/Tracing-Model-Outputs-to-the-Training-Data/index.html b/2023/11/17/Tracing-Model-Outputs-to-the-Training-Data/index.html index e9fecf11b..0423732b5 100644 --- a/2023/11/17/Tracing-Model-Outputs-to-the-Training-Data/index.html +++ b/2023/11/17/Tracing-Model-Outputs-to-the-Training-Data/index.html @@ -132,7 +132,7 @@
-
+
@@ -271,7 +271,7 @@

更新于 - + @@ -887,7 +887,7 @@

Attribution to Layers and Tokens

- 61 + 62 日志
diff --git "a/2023/11/24/vLLM\347\232\204PageAttention/index.html" "b/2023/11/24/vLLM\347\232\204PageAttention/index.html" index 445073bfb..b7b8290e1 100644 --- "a/2023/11/24/vLLM\347\232\204PageAttention/index.html" +++ "b/2023/11/24/vLLM\347\232\204PageAttention/index.html" @@ -97,7 +97,7 @@
-
+
@@ -236,7 +236,7 @@

更新于 - + @@ -383,8 +383,8 @@

策略赋值更新 kv cache

Tracing Model Outputs to the Training Data

-

@@ -460,7 +460,7 @@

策略赋值更新 kv cache

- 61 + 62 日志
diff --git "a/2023/11/26/IOT/\346\240\221\350\216\223\346\264\276/\346\240\221\350\216\223\346\264\276\346\212\230\350\205\276\350\256\260\345\275\225/index.html" "b/2023/11/26/IOT/\346\240\221\350\216\223\346\264\276/\346\240\221\350\216\223\346\264\276\346\212\230\350\205\276\350\256\260\345\275\225/index.html" index cdc9085ef..5ff655ba9 100644 --- "a/2023/11/26/IOT/\346\240\221\350\216\223\346\264\276/\346\240\221\350\216\223\346\264\276\346\212\230\350\205\276\350\256\260\345\275\225/index.html" +++ "b/2023/11/26/IOT/\346\240\221\350\216\223\346\264\276/\346\240\221\350\216\223\346\264\276\346\212\230\350\205\276\350\256\260\345\275\225/index.html" @@ -94,7 +94,7 @@
-
+
@@ -233,7 +233,7 @@

更新于 - + @@ -350,8 +350,8 @@

ics 部署

-
- 61 + 62 日志
diff --git "a/2023/11/26/IOT/\346\240\221\350\216\223\346\264\276/\346\240\221\350\216\223\346\264\276\346\216\247\345\210\266\346\211\213\346\234\272\350\207\252\345\212\250\346\213\215\347\205\247/index.html" "b/2023/11/26/IOT/\346\240\221\350\216\223\346\264\276/\346\240\221\350\216\223\346\264\276\346\216\247\345\210\266\346\211\213\346\234\272\350\207\252\345\212\250\346\213\215\347\205\247/index.html" index 75e47def6..041e1da34 100644 --- "a/2023/11/26/IOT/\346\240\221\350\216\223\346\264\276/\346\240\221\350\216\223\346\264\276\346\216\247\345\210\266\346\211\213\346\234\272\350\207\252\345\212\250\346\213\215\347\205\247/index.html" +++ "b/2023/11/26/IOT/\346\240\221\350\216\223\346\264\276/\346\240\221\350\216\223\346\264\276\346\216\247\345\210\266\346\211\213\346\234\272\350\207\252\345\212\250\346\213\215\347\205\247/index.html" @@ -93,7 +93,7 @@
-
+
@@ -232,7 +232,7 @@

更新于 - + @@ -340,12 +340,12 @@

成品

-
-
@@ -421,7 +421,7 @@

成品

- 61 + 62 日志
diff --git "a/2023/11/26/\345\276\205\346\225\264\347\220\206\350\204\232\346\234\254/index.html" "b/2023/11/26/\345\276\205\346\225\264\347\220\206\350\204\232\346\234\254/index.html" index ea8260cbb..a94ecca66 100644 --- "a/2023/11/26/\345\276\205\346\225\264\347\220\206\350\204\232\346\234\254/index.html" +++ "b/2023/11/26/\345\276\205\346\225\264\347\220\206\350\204\232\346\234\254/index.html" @@ -90,7 +90,7 @@
-
+
@@ -229,7 +229,7 @@

更新于 - + @@ -313,12 +313,12 @@

kill Python 相关进程

-
-
@@ -394,7 +394,7 @@

kill Python 相关进程

- 61 + 62 日志
diff --git a/2023/12/03/PyTorch/pytorch_cpp_extension/index.html b/2023/12/03/PyTorch/pytorch_cpp_extension/index.html index e374f54ad..ed9c8cadd 100644 --- a/2023/12/03/PyTorch/pytorch_cpp_extension/index.html +++ b/2023/12/03/PyTorch/pytorch_cpp_extension/index.html @@ -31,7 +31,7 @@ - + @@ -96,7 +96,7 @@
-
+
@@ -235,7 +235,7 @@

更新于 - + @@ -373,181 +373,10 @@

Tiled Matmul

  • 越大的 block size表示拆分后的矩阵个数变少，这样访问全局内存的次数会更多。但是，每个矩阵比较大，这样线程的并行度会更高。即，IO变小，计算变快。
  • 越小的 block size表示拆分后的矩阵个数变多，这样访问全局内存的次数会更少。但是，每个矩阵比较小，这样线程的并行度会更低。即，IO变大，计算变慢。
    • -

    算子融合

    -

    算子融合是将多个计算操作合并为一个计算操作，以减少计算量和内存访问次数，从而提高计算效率。比如，矩阵乘法和 softmax 的融合，可以减少一次内存访问。

    -

    softmax 的改造

    -

    softmax 的计算公式为：

    -

    \[ -\text{softmax}(x_i) = \frac{e^{x_i}}{\sum_{j=1}^n e^{x_j}} -\]

    -

    而通常为了数值稳定性（避免溢出），会先计算最大值，再减去最大值，最后再计算 softmax。即：

    -

    \[ -\text{softmax}(x_i) = \frac{e^{x_i - \max(x)}}{\sum_{j=1}^n e^{x_j - \max(x)}} -\]

    -

    在神经网络中，其值通常为一个矩阵（不只是一个维度），所以需要对每一行进行 softmax。示例代码如下：

    -
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    import torch

    X = torch.rand(4, 4)

    def my_softmax(X, dim=1):
        X -= torch.max(X, dim=dim, keepdim=True)[0]
        return torch.exp(X) / torch.sum(torch.exp(X), dim=dim, keepdim=True)

    print(torch.softmax(X, dim=1))
    print(my_softmax(X, dim=1))
    assert torch.allclose(torch.softmax(X, dim=1), my_softmax(X, dim=1))
    -

    但这样跟矩阵乘法进行算子融合是没有优势的，两个函数还是独立计算的。算子融合是需要糅合两种计算操作。所以，需要把 softmax 函数放到矩阵乘法中进行计算。即将计算过程变成一个可迭代的过程，换而言之，随着元素的增加不断更新 softmax 的结果。

    -

    Online normalizer calculation for softmax 提供了这么一种方式。简单起见，从一个元素开始介绍。现在有一个列表，里面只有一个元素 [1]，其 softmax 计算过程为：

    -
      -
    1. exp: 计算 \(e^{1}\)，得到 \([e^{1}]\)
      2. -
    2. max: 计算 max(1)，得到 1
      3. -
    3. -max: 计算 $e^{1 - 1} $，得到 \([e^{1 - 1} ]\)
      4. -
    4. sum: 计算 $e^{1- 1} $ 的和，得到 $e^{1- 1} $
      5. -
    5. softmax: 计算 \([e^{1- 1} ] / (e^{1- 1} )\)，得到 \([1]\)
      6. -
    -

    现在增加一个元素，观察有哪些变化。假设增加的元素为 2，现在有一个列表，里面有两个元素 [1, 2]，新增的 softmax 计算过程为：

    -
      -
    1. exp: 计算 \(e^{2}\)，得到 \(e^{2}\)
      2. -
    2. max: 与原来的 max(1)=1 比较，得到 max(1, 2)，得到 2
      3. -
    3. -max：所有元素更新一遍，均减去最新的 max(1, 2)=2，得到 \([e^{1- 2} , e^{2- 2} ]\)
      4. -
    4. sum: 重新计算一遍结果
      5. -
    5. softmax: 计算 \([e^{1- 2}, e^{2 - 2}] / (e^{1 - 2} + e^{2 - 2})\)
      6. -
    -

    对比发现，第 3 步会导致重新第 4 步计算一遍求和结果，但这个求和结果在第 5 步中作为分母是可以灵活调整的。原来是\(e^{1-1}\)，现在更新为\(e^{1-2}+e^{2-2}\)。假设原来的最大值为 old_max，新的最大值为 new_max，原来的元素为 old_v，则原来元素的 exp结果可以更新为

    -

    \[ -e^{old\_v-old\_max} \rightarrow e^{old\_v-old\_max+old\_max-new\_max} = e^{old\_v-old\_max} \times e^{old\_max-new\_max} -\]

    -

    而对于新加的元素，则可以在求和部分直接加上\(e^{2-2}\)，所以，代码如下

    -
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    import numpy as np

    nums = [1, 2]
    sum_, max_v = 0., 0.
    norm_v = 0.
    for i, num in enumerate(nums):
        old_max_v = max_v
        max_v = max(max_v, num)
        norm_v = norm_v*np.exp(old_max_v - max_v) + np.exp(num - max_v)
    softmax_nums = [np.exp(num-max_v)/ norm_v for num in nums]
    -

    对于两个元素这个公式是适用的，那么对于多个元素呢？答案自然也是同样适用的

    -

    \[ -\begin{aligned} -e^{v1-old\_max}+e^{v2-old\_max} &\rightarrow e^{v1-new\_max}+e^{v2-new\_max} \\ &= e^{v1-old\_max+old\_max-new\_max}+e^{v2-old\_max+old\_max-new\_max} \\ \\ &= e^{v1-old\_max} \times e^{old\_max-new\_max} + e^{v2-old\_max} \times e^{old\_max-new\_max} \\ &= (e^{v1-old\_max} - e^{v2-old\_max}) \times e^{old\_max-new\_max} -\end{aligned} -\]

    -

    所以，将其更新为二维矩阵的形式，就有

    -
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    14
    15
    16
    17
    18
    19
    20
    21
    import torch
    import numpy as np

    X = torch.rand(4, 4)

    def online_softmax(X):
        value = torch.zeros_like(X)
        for row in range(X.shape[0]):
            row_max = 0.0
            normalizer_term = 0.0
            for col in range(X.shape[1]):
                val = X[row, col]
                old_row_max = row_max
                row_max = max(old_row_max, val)
                normalizer_term = normalizer_term * np.exp(old_row_max - row_max) + np.exp(val - row_max)
            value[row, :] = torch.exp(X[row, :] - row_max) / normalizer_term
        return value

    print(torch.softmax(X, dim=1))
    print(online_softmax(X))
    assert torch.allclose(torch.softmax(X, dim=1), online_softmax(X))
    -

    softmax + 矩阵乘法

    -

    一般的融合操作是需要在矩阵乘法后，计算结果的 softmax。即

    -
    1
    2
    3
    4
    5
    import torch

    M, N, K = 4, 2, 4
    A1, A2 = torch.rand(size=(M, N)), torch.rand(size=(N, K))
    output = torch.softmax(A1 @ A2, dim=1)
    -

    在结合上面的内容后，将其改造为

    -
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    14
    15
    16
    17
    18
    19
    20
    21
    22
    23
    import torch
    import numpy as np

    M, N, K = 4, 2, 4
    A1, A2 = torch.rand(size=(M, N)), torch.rand(size=(N, K))

    def matmul_softmax(A1, A2):
        output = torch.zeros(size=(A1.shape[0], A2.shape[1]))
        for i in range(A1.shape[0]):
            row_max = 0.0
            normalizer_term = 0.0  
            for j in range(A2.shape[1]):
                val = output[i, j] = sum(map(lambda x: x[0] * x[1], zip(A1[i], A2[:, j])))
              
                old_row_max = row_max
                row_max = max(old_row_max, val)
                normalizer_term = normalizer_term * np.exp(old_row_max - row_max) + np.exp(val - row_max)
            output[i, :] = torch.exp(output[i, :] - row_max) / normalizer_term
        return output

    print(torch.softmax(A1 @ A2, dim=1))
    print(matmul_softmax(A1, A2))
    assert torch.allclose(torch.softmax(A1 @ A2, dim=1), matmul_softmax(A1, A2))
    -

    更进一步，结合 Tiled matmul，将其改造为：

    -
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    14
    15
    16
    17
    18
    19
    20
    21
    22
    23
    24
    25
    26
    27
    28
    29
    30
    31
    32
    33
    34
    35
    36
    37
    38
    39
    40
    41
    42
    43
    44
    45
    46
    47
    48
    49
    50
    51
    52
    53
    import torch
    import numpy as np

    M, N, K = 4, 6, 8
    A1, A2 = torch.rand(size=(M, N)), torch.rand(size=(N, K))

    def block_matmul(sub_A1, sub_A2):
        output = torch.zeros(size=(sub_A1.shape[0], sub_A2.shape[1]))
        for i in range(sub_A1.shape[0]):
            for j in range(sub_A2.shape[1]):
                for k in range(sub_A2.shape[0]):
                    output[i][j] += sub_A1[i][k] * sub_A2[k][j]            
        return output


    def tiled_matmul_softmax(A1, A2):
        block_size_M, block_size_N, block_size_K = 2, 3, 4
        block_M, block_N, block_K = M // block_size_M, N // block_size_N, K // block_size_K

        output = torch.zeros(size=(A1.shape[0], A2.shape[1]))
        for i in range(0, A1.shape[0], block_M):
            start_i, end_i = i, i + block_M
            row_max = torch.tensor([[0. for _ in range(block_N)] for _ in range(block_M)])
            old_row_max = torch.tensor([[0. for _ in range(block_N)] for _ in range(block_M)])
            normalizer_term = torch.tensor([[0. for _ in range(block_N)] for _ in range(block_M)])

            for j in range(0, A2.shape[1], block_N):
                start_j, end_j = j, j + block_N
                for k in range(0, A2.shape[0], block_K):
                    start_k, end_k = k, k + block_K
                    sub_A1 = A1[start_i:end_i, start_k:end_k]
                    sub_A2 = A2[start_k:end_k, start_j:end_j]
                    output[start_i:end_i, start_j:end_j] += block_matmul(sub_A1, sub_A2)

                # 这里算完了每个block的结果，所以需要将其拆分成每个block，然后再计算softmax
                for ii, row in enumerate(range(start_i, end_i)):            
                    for jj, col in enumerate(range(start_j, end_j)):
                        val = output[row][col]
                        old_row_max[ii][jj] = row_max[ii][jj]
                        row_max[ii][jj] = max(old_row_max[ii][jj], val)
                        normalizer_term[ii][jj] = normalizer_term[ii][jj] * np.exp(old_row_max[ii][jj] - row_max[ii][jj]) + np.exp(val - row_max[ii][jj])

            for ii, row in enumerate(range(start_i, end_i)):
                row_max_v, _ = torch.max(row_max, dim=1)
                # 重算 sum, 代入公式 old_v*exp(old_max - new_max)
                sum_ = torch.sum(normalizer_term[ii] * torch.exp(row_max[ii] - row_max_v[ii]))
                output[row, :] = torch.exp(output[row, :] - row_max_v[ii]) / sum_
        return output


    print(torch.softmax(A1 @ A2, dim=1))
    print(tiled_matmul_softmax(A1, A2))
    assert torch.allclose(torch.softmax(A1 @ A2, dim=1), tiled_matmul_softmax(A1, A2))
    -

    由于这里是每次计算出来是一个 block ，所以要把 block 拆分出每个元素，计算 block 中每行每列的最大值 row_max 以及分母 normalizer_term。

    -

    最后在计算 softmax 时，要计算所有 block 的最大值 torch.max(row_max, dim=1)，还需要计算分母，这里需要考虑所有的 block，可以类比于合并计算 [1, 2], [3, 4]两个列表的normalizer_term 。已知对应的 normalizer_term \(e^{1-2}+e^{2-2}\) 和 \(e^{3-4}+e^{3-4}\)，合并后的结果应当是 \(e^{1-4}+e^{2-4}+e^{3-4}+e^{3-4}\)。将其公式化写作：

    -

    \[ -\begin{aligned} -\sum_{i=1}^{m} e^{x_i-max(x)} + \sum_{i=1}^{n} e^{y_i-max(y)} & \rightarrow \sum_{i=1}^{m} e^{x_i-max(x,y)} + \sum_{i=1}^{n} e^{y_i-max(x,y)} \\ &= \sum_{i=1}^{m} e^{x_i-max(x)}e^{max(x)-max(x,y)} + \sum_{i=1}^{n} e^{y_i-max(x,y)}e^{max(y)-max(x,y)} -\end{aligned} -\]

    -

    假设 max(x,y)=max(x)，那么对应项将乘 1，这并不会对结果有任何影响。对应的代码为：

    -
    1
    sum_ = torch.sum(normalizer_term[ii] * torch.exp(row_max[ii] - row_max_v[ii]))
    -

    至此，完成了矩阵乘法和 softmax 的融合。接下来，会实现 cuda 版本以对比性能。

    -

    Cuda 实现

    -

    首先，接着之前矩阵乘法的cuda实现，

    -
    1
    2
    3
    4
    5
    6
    7
    if (row < input1.size(0) && col < input2.size(1)) {
        scalar_t value = 0.0;
        for (int k = 0; k < input1.size(1); ++k) {
            value += input1[row][k] * input2[k][col];
        }
        output[row][col] = value
    }
    -

    在循环之后，已经计算完了输出矩阵中的一项。需要继续算每行的 row_max和 normalizer_term。但由于这里是 cuda 中的某个 block，所以需要借助共享内存来通信每行的结果。

    -
    1
    2
    3
    4
    5
    // 使用共享内存，计算每个 row 的最大值
    __shared__ scalar_t row_max[16][16];
    __shared__ scalar_t normalizer_term[16][16];
    row_max[threadIdx.y][threadIdx.x] = value; // 先把计算结果放到 row_max 中，以便于比较大小
    __syncthreads(); // 这行代码是为了保证每个线程都已经计算完了，才能进行下一步的操作
    -

    计算过程分为三个步骤：1. 找到每行的最大值

    -
    1
    2
    3
    4
    5
    6
    for (int i = blockDim.x / 2; i > 0; i /= 2) {
        if (threadIdx.x < i) {
            row_max[threadIdx.y][threadIdx.x] = max(row_max[threadIdx.y][threadIdx.x], row_max[threadIdx.y][threadIdx.x + i]);
        }
        __syncthreads();
    }
    -
      -
    1. 计算每行的 softmax 的分母每项组成
      2. -
    -
    1
    2
    normalizer_term[threadIdx.y][threadIdx.x] = exp(value - row_max[threadIdx.y][0]);
    __syncthreads();
    -
      -
    1. 计算每行的 softmax 的每项之后
      2. -
    -
    1
    2
    3
    4
    5
    6
    7
    8
    for (int i = blockDim.x / 2; i > 0; i /= 2) {
        if (threadIdx.x < i) {
            normalizer_term[threadIdx.y][threadIdx.x] += normalizer_term[threadIdx.y][threadIdx.x + i];
        }
        __syncthreads();
    }
    // 最后将其更新到输出矩阵中
    output[row][col] = exp(value - row_max[threadIdx.y][0]) / normalizer_term[threadIdx.y][0];
    -

    完整代码如下（这里没有实现完整的attention，仅仅是一个矩阵乘法 + softmax 计算）：

    -
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    14
    15
    16
    17
    18
    19
    20
    21
    22
    23
    24
    25
    26
    27
    28
    29
    30
    31
    32
    33
    34
    35
    36
    37
    38
    39
    40
    41
    42
    43
    44
    45
    46
    47
    48
    49
    50
    51
    52
    53
    54
    55
    56
    57
    58
    59
    60
    61
    62
    63
    64
    65
    66
    67
    68
    69
    70
    71
    72
    73
    74
    75
    76
    77
    78
    79
    80
    81
    82
    #include <torch/extension.h>

    #include <cuda.h>
    #include <cuda_runtime.h>

    #include <vector>

    // Matrix multiply kernel
    template <typename scalar_t>
    __global__ void matrix_multiply_kernel(const torch::PackedTensorAccessor<scalar_t,2,torch::RestrictPtrTraits,size_t> input1,
                                           const torch::PackedTensorAccessor<scalar_t,2,torch::RestrictPtrTraits,size_t> input2,
                                           torch::PackedTensorAccessor<scalar_t,2,torch::RestrictPtrTraits,size_t> output) {
        int row = blockIdx.y * blockDim.y + threadIdx.y;
        int col = blockIdx.x * blockDim.x + threadIdx.x;

        if (row < input1.size(0) && col < input2.size(1)) {
            scalar_t value = 0.0;
            for (int k = 0; k < input1.size(1); ++k) {
                value += input1[row][k] * input2[k][col];
            }

            // 使用共享内存，计算每个 row 的最大值
            __shared__ scalar_t row_max[16][16];
            __shared__ scalar_t normalizer_term[16][16];
            row_max[threadIdx.y][threadIdx.x] = value;
            __syncthreads();

            for (int i = blockDim.x / 2; i > 0; i /= 2) {
                if (threadIdx.x < i) {
                    row_max[threadIdx.y][threadIdx.x] = max(row_max[threadIdx.y][threadIdx.x], row_max[threadIdx.y][threadIdx.x + i]);
                }
                __syncthreads();
            }
            // 计算每个 row 的 softmax 的分母
            normalizer_term[threadIdx.y][threadIdx.x] = exp(value - row_max[threadIdx.y][0]);

            __syncthreads();
            // 计算每个 row  normalizer_term之和
            for (int i = blockDim.x / 2; i > 0; i /= 2) {
                if (threadIdx.x < i) {
                    normalizer_term[threadIdx.y][threadIdx.x] += normalizer_term[threadIdx.y][threadIdx.x + i];
                }
                __syncthreads();
            }

            // 计算每个 row 的 softmax
            output[row][col] = exp(value - row_max[threadIdx.y][0]) / normalizer_term[threadIdx.y][0];
        }
    }

    torch::Tensor matrix_multiply(torch::Tensor input1, torch::Tensor input2) {
        int rows1 = input1.size(0);
        int cols1 = input1.size(1);
        int cols2 = input2.size(1);

        auto options = torch::TensorOptions().device(input1.device());
        torch::Tensor output = torch::zeros({rows1, cols2}, options);

        const dim3 threads(16, 16);
        const dim3 blocks((cols2 + threads.x - 1) / threads.x,
                          (rows1 + threads.y - 1) / threads.y);

        AT_DISPATCH_FLOATING_TYPES(input1.scalar_type(), "matrix_multiply_kernel", ([&] {
            matrix_multiply_kernel<<<blocks, threads>>>(
                input1.packed_accessor<scalar_t,2,torch::RestrictPtrTraits,size_t>(),
                input2.packed_accessor<scalar_t,2,torch::RestrictPtrTraits,size_t>(),
                output.packed_accessor<scalar_t,2,torch::RestrictPtrTraits,size_t>());
        }));

        return output;
    }


    // 这里偷懒没有换名字
    std::vector<torch::Tensor> attention_cuda_forward(
        torch::Tensor q,
        torch::Tensor k,
        torch::Tensor v) {

        torch::Tensor scores = matrix_multiply(q, k);
        return {scores};
    }
    -

    测试代码如下

    -
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    14
    15
    16
    17
    18
    19
    20
    21
    22
    23
    24
    25
    26
    27
    28
    29
    30
    31
    32
    33
    34
    35
    36
    37
    38
    39
    40
    import torch

    import mulsoftmax
    import timeit

    seed = 42
    torch.manual_seed(seed)
    torch.cuda.manual_seed(seed)

    def my_py_softmax(x, dim):
        e = torch.exp(x)
        s = torch.sum(e, dim=dim, keepdim=True)
        return e / s

    def py_mulsoft(q, k, v):
        # print(q@k.T)
        return torch.softmax(q @ k.T, dim=1)

    def check_forward(q, k, v):
        baseline_values = py_mulsoft(q, k, v)
        cpp_values = mulsoftmax.forward(q, k, v)[0]

        print("base o", baseline_values)
        print("cpp  o", cpp_values)
        print(torch.all(torch.isclose(baseline_values, cpp_values)))

    def compare_time(loop=100):
        q, k, v = torch.rand(size=(m, n), device=device), torch.rand(size=(m, n), device=device), torch.rand(size=(m, n), device=device)
        print("py", timeit.timeit(lambda: py_mulsoft(q, k, v), number=loop))
        print("cpp", timeit.timeit(lambda: mulsoftmax.forward(q, k, v)[0], number=loop))

    if __name__ == "__main__":
        m, n = 16, 40
        device = "cuda"
        q, k, v = torch.rand(size=(m, n), device=device), torch.rand(size=(m, n), device=device), torch.rand(size=(m, n), device=device)
        # 先检查结果是否正确
        check_forward(q, k, v)
        q, k, v = torch.rand(size=(m, n)), torch.rand(size=(m, n)), torch.rand(size=(m, n))
        # 循环1w次，对比性能差距
        compare_time(10000)
    -

    输出结果如下：

    - - - - - - - - - - - - - - - - - - - - - - - - - -
    pycuda
    00.51360.2909
    10.61430.3322
    20.73000.3608
    -

    其他实现

    -

    二维矩阵转一维矩阵，实现矩阵乘法

    -

    在某些情况下，为了降低空间复杂度，会把二维矩阵展开为一维矩阵，再进行矩阵乘法。这里在这个基础上，再实现了 softmax 的计算。其本质上是把二维矩阵转为一维矩阵，再进行矩阵乘法。原来的input1[row][k] -> input1[row * K + k]，input2[k][col] -> input2[k * N + col]。但这里暂时不好实现 softmax 融合，因为 softmax 需要计算每行的最大值，这里的一维矩阵无法直接计算每行的最大值。所以这里偷懒先使用判断的方法，如果是最后一列，则计算 softmax。因此，导致算的速度会比较慢。

    -
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    14
    15
    16
    17
    18
    19
    20
    21
    22
    23
    24
    25
    26
    27
    28
    29
    30
    31
    32
    33
    34
    35
    36
    37
    38
    39
    40
    41
    42
    43
    44
    45
    46
    47
    48
    49
    50
    51
    52
    53
    54
    55
    56
    57
    58
    59
    60
    61
    62
    63
    64
    65
    66
    67
    68
    69
    70
    71
    72
    73
    #include <torch/extension.h>
    #include <cuda.h>
    #include <cuda_runtime.h>

    #include <vector>

    #DEFINE BLOCK_SIZE 256;

    template <typename scalar_t>
    __global__ void matrix_multiply_vector_kernel(const torch::PackedTensorAccessor<scalar_t,1,torch::RestrictPtrTraits,size_t> input1,
                                           const torch::PackedTensorAccessor<scalar_t,1,torch::RestrictPtrTraits,size_t> input2,
                                           torch::PackedTensorAccessor<scalar_t,1,torch::RestrictPtrTraits,size_t> output,
                                           const int M, const int N, const int K
                                           ) {
        int index = blockIdx.x * blockDim.x + threadIdx.x;
        int row = index / N;
        int col = index % N;

        if (row < M && col < N) {
            float value = 0.0;
            for (int k = 0; k < K; ++k) {
                value += input1[row * K + k] * input2[k * N + col];
            }
            output[row * N + col] = value;
            if (col == N - 1) {
                float row_max = 0.0;
                float normalizer_term = 0.0;        
                float old_row_max = 0.0;
                for (int i = 0; i < N; ++i) {
                    old_row_max = row_max;
                    row_max = max(row_max, output[row * N + i]);
                    normalizer_term = normalizer_term * exp(old_row_max - row_max) + exp(output[row * N + i] - row_max);
                }
                for (int i = 0; i < N; ++i) {
                    output[row * N + i] = exp(output[row * N + i] - row_max) / normalizer_term;
                }
            }
        }
    }

    std::vector<torch::Tensor> matmul_vector(torch::Tensor input1, torch::Tensor input2) {
        int M = input1.size(0);
        int K = input1.size(1);
        int N = input2.size(1);

        auto options = torch::TensorOptions().device(input1.device());
        
        const dim3 threads(BLOCK_SIZE);
        const dim3 blocks((M * N + threads.x - 1) / threads.x);

        // Reshape input tensors to vectors
        auto input1_vector = input1.reshape({-1});
        auto input2_vector = input2.reshape({-1});    
        torch::Tensor output_vector = torch::zeros({M * N}, options);

        AT_DISPATCH_FLOATING_TYPES(input1_vector.scalar_type(), "matrix_multiply_vector_kernel", ([&] {
            matrix_multiply_vector_kernel<<<blocks, threads>>>(
                input1_vector.packed_accessor<scalar_t,1,torch::RestrictPtrTraits,size_t>(),
                input2_vector.packed_accessor<scalar_t,1,torch::RestrictPtrTraits,size_t>(),
                output_vector.packed_accessor<scalar_t,1,torch::RestrictPtrTraits,size_t>(),
                M, N, K
            );
        }));
        return {output_vector.reshape({M, N}), output_vector.reshape({M, N})};
    }

    std::vector<torch::Tensor> attention_cuda_forward(
        torch::Tensor q,
        torch::Tensor k,
        torch::Tensor v) {

        return matmul_vector(q, k);
    }
    -

    输出结果如下：

    - - - - - - - - - - - - - - - - - - - - - - - - - -
    pycuda
    00.42390.4907
    10.50690.4388
    20.54620.5799
    -

    尽管用了一种比较笨的方法，但是速度实际上与原生的 pytorch 相差无几。

    -

    单独实现 softmax

    -

    相较于非算子融合的写法，这里实现了一维矩阵的 softmax。相当于在计算矩阵乘法后再算 softmax。比较朴素的写法。。。

    -
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    14
    15
    16
    17
    18
    19
    20
    21
    22
    23
    24
    25
    26
    27
    28
    29
    30
    31
    32
    33
    34
    35
    36
    37
    38
    39
    40
    41
    42
    43
    44
    45
    46
    47
    48
    49
    50
    51
    52
    53
    54
    55
    56
    57
    58
    59
    template <typename scalar_t>
    __global__ void softmax_kernel(torch::PackedTensorAccessor<scalar_t,1,torch::RestrictPtrTraits,size_t> output,
                                     const int M, const int N) {
        int index = blockIdx.x * blockDim.x + threadIdx.x;
        int row = index / N;

        if (row < M) {
            float row_max = 0.0;
            float normalizer_term = 0.0;        
            float old_row_max = 0.0;

            for (int i = 0; i < N; ++i) {
                old_row_max = row_max;
                row_max = max(row_max, output[row * N + i]);
                normalizer_term = normalizer_term * exp(old_row_max - row_max) + exp(output[row * N + i] - row_max);
            }

            for (int i = 0; i < N; ++i) {
                output[row * N + i] = exp(output[row * N + i] - row_max) / normalizer_term;
            }
        }
    }

    torch::Tensor matrix_softmax_vector_softmax(torch::Tensor input1, torch::Tensor input2) {
        int M = input1.size(0);
        int K = input1.size(1);
        int N = input2.size(1);

        auto options = torch::TensorOptions().device(input1.device());
        
        const dim3 threads(BLOCK_SIZE_VECTOR);
        const dim3 blocks((M * N + threads.x - 1) / threads.x);

        // Reshape input tensors to vectors
        auto input1_vector = input1.reshape({-1});
        auto input2_vector = input2.reshape({-1});    
        torch::Tensor output_vector = torch::zeros({M * N}, options);

        // 普通一维的矩阵乘法
        AT_DISPATCH_FLOATING_TYPES(input1_vector.scalar_type(), "matrix_multiply_vector_softmax_kernel", ([&] {
            matrix_multiply_vector_softmax_kernel<<<blocks, threads>>>(
                input1_vector.packed_accessor<scalar_t,1,torch::RestrictPtrTraits,size_t>(),
                input2_vector.packed_accessor<scalar_t,1,torch::RestrictPtrTraits,size_t>(),
                output_vector.packed_accessor<scalar_t,1,torch::RestrictPtrTraits,size_t>(),
                M, N, K
            );
        }));

        cudaDeviceSynchronize();

        AT_DISPATCH_FLOATING_TYPES(output_vector.scalar_type(), "softmax_kernel", ([&] {
            softmax_kernel<<<blocks, threads>>>(
                output_vector.packed_accessor<scalar_t,1,torch::RestrictPtrTraits,size_t>(),
                M, N
            );
        }));
        return output_vector.reshape({M, N});

    }
    -

    输出结果如下：

    - - - - - - - - - - - - - - - - - - - - - - - - - -
    pycuda
    00.42391.1874
    10.50691.1101
    20.54621.3238
    -

    对比发现，这里的cuda还是会比pytorch的慢一些，这是因为没有用自带的softmax，且没有使用算子融合。所以就算速度会更慢。

    +

    下一篇：以Attention为例的算子融合

    参考

    [1] https://pytorch.org/tutorials/advanced/cpp_extension.html。

    [2] Tiled matmul

    -

    [3] Online softmax

    -

    [4] Online normalizer calculation for softmax

    @@ -588,7 +417,10 @@

    参考

    -
    +
    +
    @@ -647,7 +479,7 @@

    参考

    -
    1. 为什么
    2. 怎么写
      1. CPU 版本
        1. Python 部分
        2. Cpp 部分
        3. 测试
        4. 矩阵乘法
        5. softmax
      2. GPU 版本
        1. Python 部分
        2. Cpp 部分
        3. Cu 部分
    3. 矩阵乘法的优化
      1. Matmul
      2. Tiled Matmul
    4. 算子融合
      1. softmax 的改造
      2. softmax + 矩阵乘法
      3. Cuda 实现
      4. 其他实现
        1. 二维矩阵转一维矩阵，实现矩阵乘法
        2. 单独实现 softmax
    5. 参考
    +
    1. 为什么
    2. 怎么写
      1. CPU 版本
        1. Python 部分
        2. Cpp 部分
        3. 测试
        4. 矩阵乘法
        5. softmax
      2. GPU 版本
        1. Python 部分
        2. Cpp 部分
        3. Cu 部分
    3. 矩阵乘法的优化
      1. Matmul
      2. Tiled Matmul
    4. 参考
    @@ -663,7 +495,7 @@

    参考

    - 61 + 62 日志
    diff --git "a/2023/12/03/Python/Python\350\207\252\345\212\250\345\257\274\345\205\245\345\214\205/index.html" "b/2023/12/03/Python/Python\350\207\252\345\212\250\345\257\274\345\205\245\345\214\205/index.html" index b05c077b0..6005efb72 100644 --- "a/2023/12/03/Python/Python\350\207\252\345\212\250\345\257\274\345\205\245\345\214\205/index.html" +++ "b/2023/12/03/Python/Python\350\207\252\345\212\250\345\257\274\345\205\245\345\214\205/index.html" @@ -90,7 +90,7 @@
    -
    +
    @@ -222,7 +222,7 @@

    发表于 - + @@ -374,7 +374,7 @@

    解决方案

    - 61 + 62 日志
    diff --git "a/2023/12/03/\346\270\251\345\272\246\345\217\257\350\247\206\345\214\226/index.html" "b/2023/12/03/\346\270\251\345\272\246\345\217\257\350\247\206\345\214\226/index.html" index 9bb1dbc0a..dfa90d689 100644 --- "a/2023/12/03/\346\270\251\345\272\246\345\217\257\350\247\206\345\214\226/index.html" +++ "b/2023/12/03/\346\270\251\345\272\246\345\217\257\350\247\206\345\214\226/index.html" @@ -91,7 +91,7 @@
    -
    +
    @@ -223,7 +223,7 @@

    发表于 - + @@ -376,7 +376,7 @@

    代码实现

    - 61 + 62 日志
    diff --git "a/2023/12/20/PyTorch/\344\273\245Attention\344\270\272\344\276\213\347\232\204\347\256\227\345\255\220\350\236\215\345\220\210/index.html" "b/2023/12/20/PyTorch/\344\273\245Attention\344\270\272\344\276\213\347\232\204\347\256\227\345\255\220\350\236\215\345\220\210/index.html" new file mode 100644 index 000000000..1586c7c7d --- /dev/null +++ "b/2023/12/20/PyTorch/\344\273\245Attention\344\270\272\344\276\213\347\232\204\347\256\227\345\255\220\350\236\215\345\220\210/index.html" @@ -0,0 +1,816 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + 以Attention为例的算子融合 | Wnma's Blogs + + + + + + + + + + + + + + + + +
    +
    + +
    +
    +
    +
    + + + +
    +
    + +
    + + + +

    Wnma's Blogs

    + +     +
    + +
    +
    + +
    +
    +
    + + + + + + + + +
    +
    +
    + + + +
    + +
    + + + +
    +
    +
    + +
    +
    + +
    +
    + +
    +
    + + +
    + + 0% +
    + + +
    +
    +
    + + +
    + + + + + +
    + + + + + +
    +

    + 以Attention为例的算子融合 +

    + +
    + + + + + 发表于 + + + + + + + + + 更新于 + + + + + + + 分类于 + + + + + + + +
    +
    + + + + +
    + + +

    以 Attention 作为例子介绍算子融合的写法

    +

    代码

    + +

    接 用 Cpp 写 PyTorch 的插件

    +

    矩阵乘法+Softmax算子融合

    +

    算子融合是将多个计算操作合并为一个计算操作，以减少计算量和内存访问次数，从而提高计算效率。比如，矩阵乘法和 softmax 的融合，可以减少一次内存访问。

    +

    Online softmax

    +

    softmax 的计算公式为：

    +

    \[ +\text{softmax}(x_i) = \frac{e^{x_i}}{\sum_{j=1}^n e^{x_j}} +\]

    +

    而通常为了数值稳定性（避免溢出），会先计算最大值，再减去最大值，最后再计算 softmax。即：

    +

    \[ +\text{softmax}(x_i) = \frac{e^{x_i - \max(x)}}{\sum_{j=1}^n e^{x_j - \max(x)}} +\]

    +

    在神经网络中，其值通常为一个矩阵（不只是一个维度），所以需要对每一行进行 softmax。示例代码如下：

    +
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    import torch

    X = torch.rand(4, 4)

    def my_softmax(X, dim=1):
        X -= torch.max(X, dim=dim, keepdim=True)[0]
        return torch.exp(X) / torch.sum(torch.exp(X), dim=dim, keepdim=True)

    print(torch.softmax(X, dim=1))
    print(my_softmax(X, dim=1))
    assert torch.allclose(torch.softmax(X, dim=1), my_softmax(X, dim=1))
    +

    但这样跟矩阵乘法进行算子融合是没有优势的，两个函数还是独立计算的。算子融合是需要糅合两种计算操作。所以，需要把 softmax 函数放到矩阵乘法中进行计算。即将计算过程变成一个可迭代的过程，换而言之，随着元素的增加不断更新 softmax 的结果。

    +

    Online normalizer calculation for softmax 提供了这么一种方式。简单起见，从一个元素开始介绍。现在有一个列表，里面只有一个元素 [1]，其 softmax 计算过程为：

    +
      +
    1. exp: 计算 \(e^{1}\)，得到 \([e^{1}]\)
      2. +
    2. max: 计算 max(1)，得到 1
      3. +
    3. -max: 计算 $e^{1 - 1} $，得到 \([e^{1 - 1} ]\)
      4. +
    4. sum: 计算 $e^{1- 1} $ 的和，得到 $e^{1- 1} $
      5. +
    5. softmax: 计算 \([e^{1- 1} ] / (e^{1- 1} )\)，得到 \([1]\)
      6. +
    +

    现在增加一个元素，观察有哪些变化。假设增加的元素为 2，现在有一个列表，里面有两个元素 [1, 2]，新增的 softmax 计算过程为：

    +
      +
    1. exp: 计算 \(e^{2}\)，得到 \(e^{2}\)
      2. +
    2. max: 与原来的 max(1)=1 比较，得到 max(1, 2)，得到 2
      3. +
    3. -max：所有元素更新一遍，均减去最新的 max(1, 2)=2，得到 \([e^{1- 2} , e^{2- 2} ]\)
      4. +
    4. sum: 重新计算一遍结果
      5. +
    5. softmax: 计算 \([e^{1- 2}, e^{2 - 2}] / (e^{1 - 2} + e^{2 - 2})\)
      6. +
    +

    对比发现，第 3 步会导致重新第 4 步计算一遍求和结果，但这个求和结果在第 5 步中作为分母是可以灵活调整的。原来是\(e^{1-1}\)，现在更新为\(e^{1-2}+e^{2-2}\)。假设原来的最大值为 old_max，新的最大值为 new_max，原来的元素为 old_v，则原来元素的 exp结果可以更新为

    +

    \[ +e^{old\_v-old\_max} \rightarrow e^{old\_v-old\_max+old\_max-new\_max} = e^{old\_v-old\_max} \times e^{old\_max-new\_max} +\]

    +

    而对于新加的元素，则可以在求和部分直接加上\(e^{2-2}\)，所以，代码如下

    +
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    import numpy as np

    nums = [1, 2]
    sum_, max_v = 0., 0.
    norm_v = 0.
    for i, num in enumerate(nums):
        old_max_v = max_v
        max_v = max(max_v, num)
        norm_v = norm_v*np.exp(old_max_v - max_v) + np.exp(num - max_v)
    softmax_nums = [np.exp(num-max_v)/ norm_v for num in nums]
    +

    对于两个元素这个公式是适用的，那么对于多个元素呢？答案自然也是同样适用的

    +

    \[ +\begin{aligned} +e^{v1-old\_max}+e^{v2-old\_max} &\rightarrow e^{v1-new\_max}+e^{v2-new\_max} \\ &= e^{v1-old\_max+old\_max-new\_max}+e^{v2-old\_max+old\_max-new\_max} \\ \\ &= e^{v1-old\_max} \times e^{old\_max-new\_max} + e^{v2-old\_max} \times e^{old\_max-new\_max} \\ &= (e^{v1-old\_max} - e^{v2-old\_max}) \times e^{old\_max-new\_max} +\end{aligned} +\]

    +

    所以，将其更新为二维矩阵的形式，就有

    +
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    14
    15
    16
    17
    18
    19
    20
    21
    import torch
    import numpy as np

    X = torch.rand(4, 4)

    def online_softmax(X):
        value = torch.zeros_like(X)
        for row in range(X.shape[0]):
            row_max = 0.0
            normalizer_term = 0.0
            for col in range(X.shape[1]):
                val = X[row, col]
                old_row_max = row_max
                row_max = max(old_row_max, val)
                normalizer_term = normalizer_term * np.exp(old_row_max - row_max) + np.exp(val - row_max)
            value[row, :] = torch.exp(X[row, :] - row_max) / normalizer_term
        return value

    print(torch.softmax(X, dim=1))
    print(online_softmax(X))
    assert torch.allclose(torch.softmax(X, dim=1), online_softmax(X))
    +

    softmax + 矩阵乘法

    +

    一般的融合操作是需要在矩阵乘法后，计算结果的 softmax。即

    +
    1
    2
    3
    4
    5
    import torch

    M, N, K = 4, 2, 4
    A1, A2 = torch.rand(size=(M, N)), torch.rand(size=(N, K))
    output = torch.softmax(A1 @ A2, dim=1)
    +

    在结合上面的内容后，将其改造为

    +
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    14
    15
    16
    17
    18
    19
    20
    21
    22
    23
    import torch
    import numpy as np

    M, N, K = 4, 2, 4
    A1, A2 = torch.rand(size=(M, N)), torch.rand(size=(N, K))

    def matmul_softmax(A1, A2):
        output = torch.zeros(size=(A1.shape[0], A2.shape[1]))
        for i in range(A1.shape[0]):
            row_max = 0.0
            normalizer_term = 0.0  
            for j in range(A2.shape[1]):
                val = output[i, j] = sum(map(lambda x: x[0] * x[1], zip(A1[i], A2[:, j])))
              
                old_row_max = row_max
                row_max = max(old_row_max, val)
                normalizer_term = normalizer_term * np.exp(old_row_max - row_max) + np.exp(val - row_max)
            output[i, :] = torch.exp(output[i, :] - row_max) / normalizer_term
        return output

    print(torch.softmax(A1 @ A2, dim=1))
    print(matmul_softmax(A1, A2))
    assert torch.allclose(torch.softmax(A1 @ A2, dim=1), matmul_softmax(A1, A2))
    +

    更进一步，结合 Tiled matmul，将其改造为：

    +
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    14
    15
    16
    17
    18
    19
    20
    21
    22
    23
    24
    25
    26
    27
    28
    29
    30
    31
    32
    33
    34
    35
    36
    37
    38
    39
    40
    41
    42
    43
    44
    45
    46
    47
    48
    49
    50
    51
    52
    53
    import torch
    import numpy as np

    M, N, K = 4, 6, 8
    A1, A2 = torch.rand(size=(M, N)), torch.rand(size=(N, K))

    def block_matmul(sub_A1, sub_A2):
        output = torch.zeros(size=(sub_A1.shape[0], sub_A2.shape[1]))
        for i in range(sub_A1.shape[0]):
            for j in range(sub_A2.shape[1]):
                for k in range(sub_A2.shape[0]):
                    output[i][j] += sub_A1[i][k] * sub_A2[k][j]            
        return output


    def tiled_matmul_softmax(A1, A2):
        block_size_M, block_size_N, block_size_K = 2, 3, 4
        block_M, block_N, block_K = M // block_size_M, N // block_size_N, K // block_size_K

        output = torch.zeros(size=(A1.shape[0], A2.shape[1]))
        for i in range(0, A1.shape[0], block_M):
            start_i, end_i = i, i + block_M
            row_max = torch.tensor([[0. for _ in range(block_N)] for _ in range(block_M)])
            old_row_max = torch.tensor([[0. for _ in range(block_N)] for _ in range(block_M)])
            normalizer_term = torch.tensor([[0. for _ in range(block_N)] for _ in range(block_M)])

            for j in range(0, A2.shape[1], block_N):
                start_j, end_j = j, j + block_N
                for k in range(0, A2.shape[0], block_K):
                    start_k, end_k = k, k + block_K
                    sub_A1 = A1[start_i:end_i, start_k:end_k]
                    sub_A2 = A2[start_k:end_k, start_j:end_j]
                    output[start_i:end_i, start_j:end_j] += block_matmul(sub_A1, sub_A2)

                # 这里算完了每个block的结果，所以需要将其拆分成每个block，然后再计算softmax
                for ii, row in enumerate(range(start_i, end_i)):            
                    for jj, col in enumerate(range(start_j, end_j)):
                        val = output[row][col]
                        old_row_max[ii][jj] = row_max[ii][jj]
                        row_max[ii][jj] = max(old_row_max[ii][jj], val)
                        normalizer_term[ii][jj] = normalizer_term[ii][jj] * np.exp(old_row_max[ii][jj] - row_max[ii][jj]) + np.exp(val - row_max[ii][jj])

            for ii, row in enumerate(range(start_i, end_i)):
                row_max_v, _ = torch.max(row_max, dim=1)
                # 重算 sum, 代入公式 old_v*exp(old_max - new_max)
                sum_ = torch.sum(normalizer_term[ii] * torch.exp(row_max[ii] - row_max_v[ii]))
                output[row, :] = torch.exp(output[row, :] - row_max_v[ii]) / sum_
        return output


    print(torch.softmax(A1 @ A2, dim=1))
    print(tiled_matmul_softmax(A1, A2))
    assert torch.allclose(torch.softmax(A1 @ A2, dim=1), tiled_matmul_softmax(A1, A2))
    +

    由于这里是每次计算出来是一个 block ，所以要把 block 拆分出每个元素，计算 block 中每行每列的最大值 row_max 以及分母 normalizer_term。

    +

    最后在计算 softmax 时，要计算所有 block 的最大值 torch.max(row_max, dim=1)，还需要计算分母，这里需要考虑所有的 block，可以类比于合并计算 [1, 2], [3, 4]两个列表的normalizer_term 。已知对应的 normalizer_term \(e^{1-2}+e^{2-2}\) 和 \(e^{3-4}+e^{3-4}\)，合并后的结果应当是 \(e^{1-4}+e^{2-4}+e^{3-4}+e^{3-4}\)。将其公式化写作：

    +

    \[ +\begin{aligned} +\sum_{i=1}^{m} e^{x_i-max(x)} + \sum_{i=1}^{n} e^{y_i-max(y)} & \rightarrow \sum_{i=1}^{m} e^{x_i-max(x,y)} + \sum_{i=1}^{n} e^{y_i-max(x,y)} \\ &= \sum_{i=1}^{m} e^{x_i-max(x)}e^{max(x)-max(x,y)} + \sum_{i=1}^{n} e^{y_i-max(x,y)}e^{max(y)-max(x,y)} +\end{aligned} +\]

    +

    假设 max(x,y)=max(x)，那么对应项将乘 1，这并不会对结果有任何影响。对应的代码为：

    +
    1
    sum_ = torch.sum(normalizer_term[ii] * torch.exp(row_max[ii] - row_max_v[ii]))
    +

    至此，完成了矩阵乘法和 softmax 的融合。接下来，会实现 cuda 版本以对比性能。

    +

    CUDA 实现

    +

    首先，接着之前矩阵乘法的cuda实现，

    +
    1
    2
    3
    4
    5
    6
    7
    if (row < input1.size(0) && col < input2.size(1)) {
        scalar_t value = 0.0;
        for (int k = 0; k < input1.size(1); ++k) {
            value += input1[row][k] * input2[k][col];
        }
        output[row][col] = value
    }
    +

    在循环之后，已经计算完了输出矩阵中的一项。需要继续算每行的 row_max和 normalizer_term。但由于这里是 cuda 中的某个 block，所以需要借助共享内存来通信每行的结果。

    +
    1
    2
    3
    4
    5
    // 使用共享内存，计算每个 row 的最大值
    __shared__ scalar_t row_max[16][16];
    __shared__ scalar_t normalizer_term[16][16];
    row_max[threadIdx.y][threadIdx.x] = value; // 先把计算结果放到 row_max 中，以便于比较大小
    __syncthreads(); // 这行代码是为了保证每个线程都已经计算完了，才能进行下一步的操作
    +

    计算过程分为三个步骤：1. 找到每行的最大值

    +
    1
    2
    3
    4
    5
    6
    for (int i = blockDim.x / 2; i > 0; i /= 2) {
        if (threadIdx.x < i) {
            row_max[threadIdx.y][threadIdx.x] = max(row_max[threadIdx.y][threadIdx.x], row_max[threadIdx.y][threadIdx.x + i]);
        }
        __syncthreads();
    }
    +
      +
    1. 计算每行的 softmax 的分母每项组成
      2. +
    +
    1
    2
    normalizer_term[threadIdx.y][threadIdx.x] = exp(value - row_max[threadIdx.y][0]);
    __syncthreads();
    +
      +
    1. 计算每行的 softmax 的每项之后
      2. +
    +
    1
    2
    3
    4
    5
    6
    7
    8
    for (int i = blockDim.x / 2; i > 0; i /= 2) {
        if (threadIdx.x < i) {
            normalizer_term[threadIdx.y][threadIdx.x] += normalizer_term[threadIdx.y][threadIdx.x + i];
        }
        __syncthreads();
    }
    // 最后将其更新到输出矩阵中
    output[row][col] = exp(value - row_max[threadIdx.y][0]) / normalizer_term[threadIdx.y][0];
    +

    完整代码如下（这里没有实现完整的attention，仅仅是一个矩阵乘法 + softmax 计算）：

    +
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    14
    15
    16
    17
    18
    19
    20
    21
    22
    23
    24
    25
    26
    27
    28
    29
    30
    31
    32
    33
    34
    35
    36
    37
    38
    39
    40
    41
    42
    43
    44
    45
    46
    47
    48
    49
    50
    51
    52
    53
    54
    55
    56
    57
    58
    59
    60
    61
    62
    63
    64
    65
    66
    67
    68
    69
    70
    71
    72
    73
    74
    75
    76
    77
    78
    79
    80
    81
    82
    #include <torch/extension.h>

    #include <cuda.h>
    #include <cuda_runtime.h>

    #include <vector>

    // Matrix multiply kernel
    template <typename scalar_t>
    __global__ void matrix_multiply_kernel(const torch::PackedTensorAccessor<scalar_t,2,torch::RestrictPtrTraits,size_t> input1,
                                           const torch::PackedTensorAccessor<scalar_t,2,torch::RestrictPtrTraits,size_t> input2,
                                           torch::PackedTensorAccessor<scalar_t,2,torch::RestrictPtrTraits,size_t> output) {
        int row = blockIdx.y * blockDim.y + threadIdx.y;
        int col = blockIdx.x * blockDim.x + threadIdx.x;

        if (row < input1.size(0) && col < input2.size(1)) {
            scalar_t value = 0.0;
            for (int k = 0; k < input1.size(1); ++k) {
                value += input1[row][k] * input2[k][col];
            }

            // 使用共享内存，计算每个 row 的最大值
            __shared__ scalar_t row_max[16][16];
            __shared__ scalar_t normalizer_term[16][16];
            row_max[threadIdx.y][threadIdx.x] = value;
            __syncthreads();

            for (int i = blockDim.x / 2; i > 0; i /= 2) {
                if (threadIdx.x < i) {
                    row_max[threadIdx.y][threadIdx.x] = max(row_max[threadIdx.y][threadIdx.x], row_max[threadIdx.y][threadIdx.x + i]);
                }
                __syncthreads();
            }
            // 计算每个 row 的 softmax 的分母
            normalizer_term[threadIdx.y][threadIdx.x] = exp(value - row_max[threadIdx.y][0]);

            __syncthreads();
            // 计算每个 row  normalizer_term之和
            for (int i = blockDim.x / 2; i > 0; i /= 2) {
                if (threadIdx.x < i) {
                    normalizer_term[threadIdx.y][threadIdx.x] += normalizer_term[threadIdx.y][threadIdx.x + i];
                }
                __syncthreads();
            }

            // 计算每个 row 的 softmax
            output[row][col] = exp(value - row_max[threadIdx.y][0]) / normalizer_term[threadIdx.y][0];
        }
    }

    torch::Tensor matrix_multiply(torch::Tensor input1, torch::Tensor input2) {
        int rows1 = input1.size(0);
        int cols1 = input1.size(1);
        int cols2 = input2.size(1);

        auto options = torch::TensorOptions().device(input1.device());
        torch::Tensor output = torch::zeros({rows1, cols2}, options);

        const dim3 threads(16, 16);
        const dim3 blocks((cols2 + threads.x - 1) / threads.x,
                          (rows1 + threads.y - 1) / threads.y);

        AT_DISPATCH_FLOATING_TYPES(input1.scalar_type(), "matrix_multiply_kernel", ([&] {
            matrix_multiply_kernel<<<blocks, threads>>>(
                input1.packed_accessor<scalar_t,2,torch::RestrictPtrTraits,size_t>(),
                input2.packed_accessor<scalar_t,2,torch::RestrictPtrTraits,size_t>(),
                output.packed_accessor<scalar_t,2,torch::RestrictPtrTraits,size_t>());
        }));

        return output;
    }


    // 这里偷懒没有换名字
    std::vector<torch::Tensor> attention_cuda_forward(
        torch::Tensor q,
        torch::Tensor k,
        torch::Tensor v) {

        torch::Tensor scores = matrix_multiply(q, k);
        return {scores};
    }
    +

    测试代码如下

    +
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    14
    15
    16
    17
    18
    19
    20
    21
    22
    23
    24
    25
    26
    27
    28
    29
    30
    31
    32
    33
    34
    35
    36
    37
    38
    39
    40
    import torch

    import mulsoftmax
    import timeit

    seed = 42
    torch.manual_seed(seed)
    torch.cuda.manual_seed(seed)

    def my_py_softmax(x, dim):
        e = torch.exp(x)
        s = torch.sum(e, dim=dim, keepdim=True)
        return e / s

    def py_mulsoft(q, k, v):
        # print(q@k.T)
        return torch.softmax(q @ k.T, dim=1)

    def check_forward(q, k, v):
        baseline_values = py_mulsoft(q, k, v)
        cpp_values = mulsoftmax.forward(q, k, v)[0]

        print("base o", baseline_values)
        print("cpp  o", cpp_values)
        print(torch.all(torch.isclose(baseline_values, cpp_values)))

    def compare_time(loop=100):
        q, k, v = torch.rand(size=(m, n), device=device), torch.rand(size=(m, n), device=device), torch.rand(size=(m, n), device=device)
        print("py", timeit.timeit(lambda: py_mulsoft(q, k, v), number=loop))
        print("cpp", timeit.timeit(lambda: mulsoftmax.forward(q, k, v)[0], number=loop))

    if __name__ == "__main__":
        m, n = 16, 40
        device = "cuda"
        q, k, v = torch.rand(size=(m, n), device=device), torch.rand(size=(m, n), device=device), torch.rand(size=(m, n), device=device)
        # 先检查结果是否正确
        check_forward(q, k, v)
        q, k, v = torch.rand(size=(m, n)), torch.rand(size=(m, n)), torch.rand(size=(m, n))
        # 循环1w次，对比性能差距
        compare_time(10000)
    +

    输出结果如下：

    + + + + + + + + + + + + + + + + + + + + + + + + + +
    pycuda
    00.51360.2909
    10.61430.3322
    20.73000.3608
    +

    其他实现

    +

    二维矩阵转一维矩阵，实现矩阵乘法

    +

    在某些情况下，为了降低空间复杂度，会把二维矩阵展开为一维矩阵，再进行矩阵乘法。这里在这个基础上，再实现了 softmax 的计算。其本质上是把二维矩阵转为一维矩阵，再进行矩阵乘法。原来的input1[row][k] -> input1[row * K + k]，input2[k][col] -> input2[k * N + col]。但这里暂时不好实现 softmax 融合，因为 softmax 需要计算每行的最大值，这里的一维矩阵无法直接计算每行的最大值。所以这里偷懒先使用判断的方法，如果是最后一列，则计算 softmax。因此，导致算的速度会比较慢。

    +
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    14
    15
    16
    17
    18
    19
    20
    21
    22
    23
    24
    25
    26
    27
    28
    29
    30
    31
    32
    33
    34
    35
    36
    37
    38
    39
    40
    41
    42
    43
    44
    45
    46
    47
    48
    49
    50
    51
    52
    53
    54
    55
    56
    57
    58
    59
    60
    61
    62
    63
    64
    65
    66
    67
    68
    69
    70
    71
    72
    73
    #include <torch/extension.h>
    #include <cuda.h>
    #include <cuda_runtime.h>

    #include <vector>

    #DEFINE BLOCK_SIZE 256;

    template <typename scalar_t>
    __global__ void matrix_multiply_vector_kernel(const torch::PackedTensorAccessor<scalar_t,1,torch::RestrictPtrTraits,size_t> input1,
                                           const torch::PackedTensorAccessor<scalar_t,1,torch::RestrictPtrTraits,size_t> input2,
                                           torch::PackedTensorAccessor<scalar_t,1,torch::RestrictPtrTraits,size_t> output,
                                           const int M, const int N, const int K
                                           ) {
        int index = blockIdx.x * blockDim.x + threadIdx.x;
        int row = index / N;
        int col = index % N;

        if (row < M && col < N) {
            float value = 0.0;
            for (int k = 0; k < K; ++k) {
                value += input1[row * K + k] * input2[k * N + col];
            }
            output[row * N + col] = value;
            if (col == N - 1) {
                float row_max = 0.0;
                float normalizer_term = 0.0;        
                float old_row_max = 0.0;
                for (int i = 0; i < N; ++i) {
                    old_row_max = row_max;
                    row_max = max(row_max, output[row * N + i]);
                    normalizer_term = normalizer_term * exp(old_row_max - row_max) + exp(output[row * N + i] - row_max);
                }
                for (int i = 0; i < N; ++i) {
                    output[row * N + i] = exp(output[row * N + i] - row_max) / normalizer_term;
                }
            }
        }
    }

    std::vector<torch::Tensor> matmul_vector(torch::Tensor input1, torch::Tensor input2) {
        int M = input1.size(0);
        int K = input1.size(1);
        int N = input2.size(1);

        auto options = torch::TensorOptions().device(input1.device());
        
        const dim3 threads(BLOCK_SIZE);
        const dim3 blocks((M * N + threads.x - 1) / threads.x);

        // Reshape input tensors to vectors
        auto input1_vector = input1.reshape({-1});
        auto input2_vector = input2.reshape({-1});    
        torch::Tensor output_vector = torch::zeros({M * N}, options);

        AT_DISPATCH_FLOATING_TYPES(input1_vector.scalar_type(), "matrix_multiply_vector_kernel", ([&] {
            matrix_multiply_vector_kernel<<<blocks, threads>>>(
                input1_vector.packed_accessor<scalar_t,1,torch::RestrictPtrTraits,size_t>(),
                input2_vector.packed_accessor<scalar_t,1,torch::RestrictPtrTraits,size_t>(),
                output_vector.packed_accessor<scalar_t,1,torch::RestrictPtrTraits,size_t>(),
                M, N, K
            );
        }));
        return {output_vector.reshape({M, N}), output_vector.reshape({M, N})};
    }

    std::vector<torch::Tensor> attention_cuda_forward(
        torch::Tensor q,
        torch::Tensor k,
        torch::Tensor v) {

        return matmul_vector(q, k);
    }
    +

    输出结果如下：

    + + + + + + + + + + + + + + + + + + + + + + + + + +
    pycuda
    00.42390.4907
    10.50690.4388
    20.54620.5799
    +

    尽管用了一种比较笨的方法，但是速度实际上与原生的 pytorch 相差无几。

    +

    单独实现 softmax

    +

    相较于非算子融合的写法，这里实现了一维矩阵的 softmax。相当于在计算矩阵乘法后再算 softmax。比较朴素的写法。。。

    +
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    14
    15
    16
    17
    18
    19
    20
    21
    22
    23
    24
    25
    26
    27
    28
    29
    30
    31
    32
    33
    34
    35
    36
    37
    38
    39
    40
    41
    42
    43
    44
    45
    46
    47
    48
    49
    50
    51
    52
    53
    54
    55
    56
    57
    58
    59
    template <typename scalar_t>
    __global__ void softmax_kernel(torch::PackedTensorAccessor<scalar_t,1,torch::RestrictPtrTraits,size_t> output,
                                     const int M, const int N) {
        int index = blockIdx.x * blockDim.x + threadIdx.x;
        int row = index / N;

        if (row < M) {
            float row_max = 0.0;
            float normalizer_term = 0.0;        
            float old_row_max = 0.0;

            for (int i = 0; i < N; ++i) {
                old_row_max = row_max;
                row_max = max(row_max, output[row * N + i]);
                normalizer_term = normalizer_term * exp(old_row_max - row_max) + exp(output[row * N + i] - row_max);
            }

            for (int i = 0; i < N; ++i) {
                output[row * N + i] = exp(output[row * N + i] - row_max) / normalizer_term;
            }
        }
    }

    torch::Tensor matrix_softmax_vector_softmax(torch::Tensor input1, torch::Tensor input2) {
        int M = input1.size(0);
        int K = input1.size(1);
        int N = input2.size(1);

        auto options = torch::TensorOptions().device(input1.device());
        
        const dim3 threads(BLOCK_SIZE_VECTOR);
        const dim3 blocks((M * N + threads.x - 1) / threads.x);

        // Reshape input tensors to vectors
        auto input1_vector = input1.reshape({-1});
        auto input2_vector = input2.reshape({-1});    
        torch::Tensor output_vector = torch::zeros({M * N}, options);

        // 普通一维的矩阵乘法
        AT_DISPATCH_FLOATING_TYPES(input1_vector.scalar_type(), "matrix_multiply_vector_softmax_kernel", ([&] {
            matrix_multiply_vector_softmax_kernel<<<blocks, threads>>>(
                input1_vector.packed_accessor<scalar_t,1,torch::RestrictPtrTraits,size_t>(),
                input2_vector.packed_accessor<scalar_t,1,torch::RestrictPtrTraits,size_t>(),
                output_vector.packed_accessor<scalar_t,1,torch::RestrictPtrTraits,size_t>(),
                M, N, K
            );
        }));

        cudaDeviceSynchronize();

        AT_DISPATCH_FLOATING_TYPES(output_vector.scalar_type(), "softmax_kernel", ([&] {
            softmax_kernel<<<blocks, threads>>>(
                output_vector.packed_accessor<scalar_t,1,torch::RestrictPtrTraits,size_t>(),
                M, N
            );
        }));
        return output_vector.reshape({M, N});

    }
    +

    输出结果如下：

    + + + + + + + + + + + + + + + + + + + + + + + + + +
    pycuda
    00.42391.1874
    10.50691.1101
    20.54621.3238
    +

    对比发现，这里的cuda还是会比pytorch的慢一些，这是因为没有用自带的softmax，且没有使用算子融合。所以就算速度会更慢。

    +

    Attention 的算子融合写法

    +

    Attention 的公式如下所示，在上面已经是实现了一次矩阵乘法 + softmax 计算的融合，接下来需要再将其与新的矩阵乘法融合。

    +
    1
    2
    def py_attention(q, k, v):
        return torch.softmax(q @ k.T, dim=1) @ v
    +

    相较于之前的先计算矩阵乘法再计算 softmax，这里是先计算 softmax 再计算矩阵乘法。重新回到开始，假设q @ v.T=out矩阵乘法结果为

    +

    \[ +[[1, 2, 3, 4], + [1, 2, 3, 4]] +\]

    +

    在计算过程中，往往是先计算得到前半部分，再得到后半部分，即

    +

    \[ +[[1, 2], + [1, 2]] +\] 和 \[ +[[3, 4], + [3, 4]] +\]

    +

    第一次计算

    +

    对于out，每行的计算过程是独立的，所以可以只模拟一行一列的情况来表示所有行的计算过程。假设k的第一列全为1。第一次计算为

    +

    \[ +[1, 2] +\]

    +
      +
    1. \(v_1 = \max(1, 2)\)
      2. +
    2. \(\text{denominator1} = \exp(1-v_1) + \exp(2-v_1)\)
      3. +
    3. \([\exp(1-v_1) / \text{denominator1}, \exp(2-v_1) / \text{denominator1}]\)
      4. +
    4. 此时，再乘上k的第一列前两项（因为这里只有两项）
      5. +
    +

    \[ +[[1], + [1]] +\]

    +
      +
    1. 有\(\exp(1-v_1) / \text{denominator1} * 1 + \exp(2-v_1) / \text{denominator1} * 1\)
      2. +
    +

    第二次计算

    +

    \[ +[3, 4] +\]

    +
      +
    1. \(v_2 = \max(3, 4, v_1)\)
      2. +
    2. \(\text{denominator2} = \text{denominator1}*(exp(v_1-v_2))+d_t\)，其中，\(d_t=\exp(3-v_2) + \exp(4-v_2)\)， 见上文的 Online Softmax
      3. +
    3. \([\exp(3-v_2) / \text{denominator2}, \exp(4-v_2) / \text{denominator2}]\)
      4. +
    4. 此时，再乘上k的第一列后两项
      5. +
    +

    \[ +[[1], + [1]] +\]

    +
      +
    1. 有\(\exp(3-v2) / \text{denominator2} * 1 + \exp(4-v2) / \text{denominator2} * 1\)
      2. +
    +

    所以现在计算的结果是

    +

    \[ +\begin{aligned} + \exp(1-v_1) / \text{denominator1} * 1 + \exp(2-v_1) / \text{denominator1} * 1 \\ ++ \exp(3-v_2) / \text{denominator2} * 1 + \exp(4-v_2) / \text{denominator2} * 1 +\end{aligned} +\]

    +

    然而，正确的结果应当是

    +

    \[ +\begin{aligned} +& \exp(1-v_2) / \text{denominator2} * 1 + \exp(2-v_2) / \text{denominator2} * 1 \\ +& + \exp(3-v_2) / \text{denominator2} * 1 + \exp(4-v_2) / \text{denominator2} * 1 +\end{aligned} +\]

    +

    校准之前的结果

    +

    对比发现，只有之前的结果是有偏差的，所以只需要对之前的结果进行校准，即乘上一个系数。这个系数怎么计算呢？以第一项为例，用实际的结果除以之前的结果，有

    +

    \[ +\begin{aligned} +& \frac{\exp(1-v_1)}{\text{denominator1}}\times \frac{\text{denominator1}}{\text{denominator2}} \times \exp(v_1-v_2) = \frac{\exp(1-v_2)}{\text{denominator2}} \\ +\end{aligned} +\]

    +

    并且，对于第二项，也是一样的，所以这个系数为

    +

    \[ +\frac{\text{denominator1}*\exp(v_1-v_2)}{\text{denominator2}} +\]

    +

    因此，对于每次校准结果，只需要先乘上这个系数（当前的denominator除以上一轮的denominator），最后再计算第五步。

    +

    在计算过程中，需要存denominator这个变量，同时还需要额外存一次之前的最大值，以进行校准。

    +

    Python 代码实现

    +

    为了便于理解，约定变量名为

    +
      +
    • old_row_max：之前的最大值
      • +
    • row_max：最新的最大值
      • +
    • denominator：之前的分母
      • +
    • mod_denominator：之前的分母进行校准后的结果
      • +
    • new_molecule：最新的分子
      • +
    • cur_denominator：分母
      • +
    +

    所以，校准公式为

    +

    \[ +\begin{aligned} +\frac{\text{denominator1}*\exp(v_1-v_2)}{\text{denominator2}} &= \frac{\text{denominator} * \exp(\text{row\_max}-\text{old\_row\_max})}{\text{cur\_denominator}} \\ +&= \frac{\text{mod\_denominator}}{\text{cur\_denominator}} +\end{aligned} +\]

    +
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    14
    15
    16
    17
    18
    19
    20
    21
    22
    23
    24
    25
    26
    27
    28
    29
    30
    31
    32
    33
    34
    35
    36
    37
    38
    39
    40
    41
    42
    43
    44
    45
    46
    47
    48
    49
    50
    51
    52
    53
    54
    55
    56
    57
    58
    59
    60
    import torch

    M = 4
    N = 6

    q = torch.rand((M, N))
    k = torch.rand((M, N))
    v = torch.rand((M, N))

    def flash_attention(q, k, v):
        output = torch.zeros(q.shape)

        block_m = 2
        block_n = 2
        block_head = N
        for i in range(0, M, block_m):
            start_i, end_i = i, i + block_m

            old_row_max = torch.zeros([block_m]) - float("inf")
            denominator = torch.zeros([block_m]) # 用于存储分母
            acc = torch.zeros([block_m, block_head])
            q_sub = q[start_i:end_i, :]
            for j in range(0, M, block_n):
                start_j, end_j = j, j + block_n
                k_sub = k[start_j:end_j, :]
                v_sub = v[start_j:end_j, :]
                qk = q_sub @ k_sub.T
                # online softmax
                row_max = torch.max(
                    torch.stack((torch.max(qk, dim=1).values, old_row_max), dim=0), dim=0
                ).values
                mod_denominator = denominator * torch.exp(old_row_max - row_max)    # 对之前的分母进行校准
                new_molecule = torch.exp(qk - row_max.reshape(-1, 1))            # 最新的分子
                cur_denominator = torch.sum(new_molecule, -1) + mod_denominator  # 对分母进行更新

                # 分子 / 分母，当前block的softmax
                new_softmax = new_molecule / torch.unsqueeze(cur_denominator, dim=1)
                # 校准系数
                acc *= torch.unsqueeze(mod_denominator / cur_denominator, dim=1)
                acc += new_softmax @ v_sub
                # 更新需要存的值
                old_row_max = row_max
                denominator = cur_denominator

            output[start_i:end_i, :] = acc

        return output


    def naive_attention(q, k, v):
        return torch.softmax(q @ k.T, dim=1) @ v


    if __name__ == "__main__":
        desired = naive_attention(q, k, v)
        actual = flash_attention(q, k, v)
        print(desired)
        print(actual)
        assert torch.allclose(desired, actual)

    +

    CUDA 实现

    +

    参考

    +

    [1] Online softmax

    +

    [2] Online normalizer calculation for softmax

    +

    [3] Flash Attention on INTEL GPU

    + +
    + + + + + + +
    + +
    + + + + +
    + + + + + + +
    + + + + +
    + + +
    + + + +
    + + + + + +
    +
    + +
    +
    + + + + +
    + + © + 2023 + + + + wnma3mz +
    +
    由 Hexo & NexT.Mist 强力驱动 +
    + + + + + + + + + + +
    +
    +
    + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/archives/1970/01/index.html b/archives/1970/01/index.html index dd29c5f97..e8cb443e5 100644 --- a/archives/1970/01/index.html +++ b/archives/1970/01/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -297,7 +297,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/1970/index.html b/archives/1970/index.html index 48fb99015..dae751cd9 100644 --- a/archives/1970/index.html +++ b/archives/1970/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -297,7 +297,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2017/04/index.html b/archives/2017/04/index.html index 5ee7481bf..9ab5cd55b 100644 --- a/archives/2017/04/index.html +++ b/archives/2017/04/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -300,7 +300,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2017/05/index.html b/archives/2017/05/index.html index ac4ad6249..9818de939 100644 --- a/archives/2017/05/index.html +++ b/archives/2017/05/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -300,7 +300,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2017/08/index.html b/archives/2017/08/index.html index 532fcc88d..f231c8b46 100644 --- a/archives/2017/08/index.html +++ b/archives/2017/08/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -340,7 +340,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2017/11/index.html b/archives/2017/11/index.html index 663283bdc..7fee8abdf 100644 --- a/archives/2017/11/index.html +++ b/archives/2017/11/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -320,7 +320,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2017/index.html b/archives/2017/index.html index 40d832e27..dbec24404 100644 --- a/archives/2017/index.html +++ b/archives/2017/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -420,7 +420,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2018/01/index.html b/archives/2018/01/index.html index 874e78cf1..ce9675af4 100644 --- a/archives/2018/01/index.html +++ b/archives/2018/01/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -320,7 +320,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2018/02/index.html b/archives/2018/02/index.html index b4177d208..482e46623 100644 --- a/archives/2018/02/index.html +++ b/archives/2018/02/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -300,7 +300,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2018/03/index.html b/archives/2018/03/index.html index c0dfd1ed2..b7cde2332 100644 --- a/archives/2018/03/index.html +++ b/archives/2018/03/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -238,8 +238,8 @@

    Wnma's Blogs

    -
    @@ -258,8 +258,8 @@

    Wnma's Blogs

    -
    @@ -360,7 +360,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2018/04/index.html b/archives/2018/04/index.html index ee2f7fbc0..527f688bb 100644 --- a/archives/2018/04/index.html +++ b/archives/2018/04/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -380,7 +380,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2018/05/index.html b/archives/2018/05/index.html index f25014997..e12af30fb 100644 --- a/archives/2018/05/index.html +++ b/archives/2018/05/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -420,7 +420,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2018/08/index.html b/archives/2018/08/index.html index 9481c6e7b..8b5d2c388 100644 --- a/archives/2018/08/index.html +++ b/archives/2018/08/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -340,7 +340,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2018/09/index.html b/archives/2018/09/index.html index bf722f4e0..b0cb8f0a2 100644 --- a/archives/2018/09/index.html +++ b/archives/2018/09/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -300,7 +300,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2018/index.html b/archives/2018/index.html index fc0e3027b..914691b62 100644 --- a/archives/2018/index.html +++ b/archives/2018/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -415,7 +415,7 @@

    Wnma's Blogs

    @@ -483,7 +483,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2018/page/2/index.html b/archives/2018/page/2/index.html index 194425b7f..01bc83fd6 100644 --- a/archives/2018/page/2/index.html +++ b/archives/2018/page/2/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -358,8 +358,8 @@

    Wnma's Blogs

    -
    @@ -378,8 +378,8 @@

    Wnma's Blogs

    -
    @@ -415,7 +415,7 @@

    Wnma's Blogs

    @@ -483,7 +483,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2018/page/3/index.html b/archives/2018/page/3/index.html index e2d7bf474..7dabad957 100644 --- a/archives/2018/page/3/index.html +++ b/archives/2018/page/3/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -275,7 +275,7 @@

    Wnma's Blogs

    @@ -343,7 +343,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2019/05/index.html b/archives/2019/05/index.html index 203dd0bcd..fe0aeda29 100644 --- a/archives/2019/05/index.html +++ b/archives/2019/05/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -320,7 +320,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2019/09/index.html b/archives/2019/09/index.html index 3f5161e6a..df0e6ee48 100644 --- a/archives/2019/09/index.html +++ b/archives/2019/09/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -300,7 +300,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2019/index.html b/archives/2019/index.html index 0d0c2b776..86f8b93d1 100644 --- a/archives/2019/index.html +++ b/archives/2019/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -340,7 +340,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2020/06/index.html b/archives/2020/06/index.html index f9cdafee9..80c6fb610 100644 --- a/archives/2020/06/index.html +++ b/archives/2020/06/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -300,7 +300,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2020/08/index.html b/archives/2020/08/index.html index 05293382e..290d37e14 100644 --- a/archives/2020/08/index.html +++ b/archives/2020/08/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -300,7 +300,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2020/09/index.html b/archives/2020/09/index.html index ec81ac89a..e7cc01009 100644 --- a/archives/2020/09/index.html +++ b/archives/2020/09/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -320,7 +320,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2020/index.html b/archives/2020/index.html index ae8cf6b29..7896ea582 100644 --- a/archives/2020/index.html +++ b/archives/2020/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -360,7 +360,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2021/02/index.html b/archives/2021/02/index.html index 7e1920e8e..14bf2d784 100644 --- a/archives/2021/02/index.html +++ b/archives/2021/02/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -320,7 +320,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2021/03/index.html b/archives/2021/03/index.html index 3ac601f9f..aa85ebe9e 100644 --- a/archives/2021/03/index.html +++ b/archives/2021/03/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -300,7 +300,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2021/07/index.html b/archives/2021/07/index.html index 58f69e6f8..4a71d47c0 100644 --- a/archives/2021/07/index.html +++ b/archives/2021/07/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -340,7 +340,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2021/10/index.html b/archives/2021/10/index.html index 2a51f6243..149653482 100644 --- a/archives/2021/10/index.html +++ b/archives/2021/10/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -300,7 +300,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2021/index.html b/archives/2021/index.html index 06ae3646a..3fe9284ac 100644 --- a/archives/2021/index.html +++ b/archives/2021/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -420,7 +420,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2023/01/index.html b/archives/2023/01/index.html index 46a209400..57e78a26f 100644 --- a/archives/2023/01/index.html +++ b/archives/2023/01/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -300,7 +300,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2023/04/index.html b/archives/2023/04/index.html index 3b0cab0b1..818618821 100644 --- a/archives/2023/04/index.html +++ b/archives/2023/04/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -300,7 +300,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2023/08/index.html b/archives/2023/08/index.html index 44af6a8e3..6369dee80 100644 --- a/archives/2023/08/index.html +++ b/archives/2023/08/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -360,7 +360,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2023/09/index.html b/archives/2023/09/index.html index ccbc3c980..02611bcf1 100644 --- a/archives/2023/09/index.html +++ b/archives/2023/09/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -300,7 +300,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2023/11/index.html b/archives/2023/11/index.html index 60a6178b1..4240f9409 100644 --- a/archives/2023/11/index.html +++ b/archives/2023/11/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -238,8 +238,8 @@

    Wnma's Blogs

    -
    @@ -258,8 +258,8 @@

    Wnma's Blogs

    -
    @@ -400,7 +400,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2023/12/index.html b/archives/2023/12/index.html index 43fc9c895..fbb266327 100644 --- a/archives/2023/12/index.html +++ b/archives/2023/12/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -206,6 +206,26 @@

    Wnma's Blogs

    2023
    + +
    @@ -340,7 +360,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2023/index.html b/archives/2023/index.html index aa618e5a8..f5a4416b4 100644 --- a/archives/2023/index.html +++ b/archives/2023/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -206,6 +206,26 @@

    Wnma's Blogs

    2023
    + +
    @@ -298,8 +318,8 @@

    Wnma's Blogs

    -
    @@ -318,8 +338,8 @@

    Wnma's Blogs

    -
    @@ -385,26 +405,6 @@

    Wnma's Blogs

    - -
    @@ -415,7 +415,7 @@

    Wnma's Blogs

    @@ -483,7 +483,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/2023/page/2/index.html b/archives/2023/page/2/index.html index a4b72ed11..7b2b6621c 100644 --- a/archives/2023/page/2/index.html +++ b/archives/2023/page/2/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -206,6 +206,26 @@

    Wnma's Blogs

    2023
    + +
    @@ -335,7 +355,7 @@

    Wnma's Blogs

    @@ -403,7 +423,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/index.html b/archives/index.html index e4e936701..1b7bf6f19 100644 --- a/archives/index.html +++ b/archives/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -206,6 +206,26 @@

    Wnma's Blogs

    2023
    + +
    @@ -298,8 +318,8 @@

    Wnma's Blogs

    -
    @@ -318,8 +338,8 @@

    Wnma's Blogs

    -
    @@ -385,26 +405,6 @@

    Wnma's Blogs

    - -
    @@ -415,7 +415,7 @@

    Wnma's Blogs

    @@ -483,7 +483,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/page/2/index.html b/archives/page/2/index.html index c2ac2f64c..a183df36e 100644 --- a/archives/page/2/index.html +++ b/archives/page/2/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -206,6 +206,26 @@

    Wnma's Blogs

    2023
    + +
    @@ -388,26 +408,6 @@

    Wnma's Blogs

    - -
    @@ -418,7 +418,7 @@

    Wnma's Blogs

    @@ -486,7 +486,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/page/3/index.html b/archives/page/3/index.html index 8adbb3ae1..60142259e 100644 --- a/archives/page/3/index.html +++ b/archives/page/3/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -206,6 +206,26 @@

    Wnma's Blogs

    2021
    + +
    @@ -391,26 +411,6 @@

    Wnma's Blogs

    - -
    @@ -421,7 +421,7 @@

    Wnma's Blogs

    @@ -489,7 +489,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/page/4/index.html b/archives/page/4/index.html index 8f8277dae..98853e9ff 100644 --- a/archives/page/4/index.html +++ b/archives/page/4/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,10 +198,33 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    +
    + 2019 +
    + +
    2018
    @@ -385,26 +408,6 @@

    Wnma's Blogs

    - -
    @@ -415,7 +418,7 @@

    Wnma's Blogs

    @@ -483,7 +486,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/page/5/index.html b/archives/page/5/index.html index 98da9d63b..94ce9239a 100644 --- a/archives/page/5/index.html +++ b/archives/page/5/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -206,6 +206,26 @@

    Wnma's Blogs

    2018
    + +
    @@ -358,8 +378,8 @@

    Wnma's Blogs

    -
    @@ -378,28 +398,8 @@

    Wnma's Blogs

    - -
    - -
    - - -
    -
    - -
    - -
    - -
    -
    @@ -415,7 +415,7 @@

    Wnma's Blogs

    @@ -483,7 +483,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/page/6/index.html b/archives/page/6/index.html index 310adbe97..421e77c5a 100644 --- a/archives/page/6/index.html +++ b/archives/page/6/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,7 +198,7 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    @@ -206,6 +206,26 @@

    Wnma's Blogs

    2018
    + +
    @@ -388,26 +408,6 @@

    Wnma's Blogs

    - -
    @@ -418,7 +418,7 @@

    Wnma's Blogs

    @@ -486,7 +486,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/archives/page/7/index.html b/archives/page/7/index.html index ac03a22b8..964c4e000 100644 --- a/archives/page/7/index.html +++ b/archives/page/7/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -198,10 +198,36 @@

    Wnma's Blogs

    - 好! 目前共计 61 篇日志。 继续努力。 + 好! 目前共计 62 篇日志。 继续努力。
    +
    + 2017 +
    + + +
    + 1970 +
    @@ -232,7 +258,7 @@

    Wnma's Blogs

    @@ -300,7 +326,7 @@

    Wnma's Blogs

    - 61 + 62 日志
    diff --git a/categories/Algorithm/index.html b/categories/Algorithm/index.html index bb10245f4..305f9fed0 100644 --- a/categories/Algorithm/index.html +++ b/categories/Algorithm/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -386,7 +386,7 @@

    Algorithm
    - 61 + 62 日志
    diff --git a/categories/Android/index.html b/categories/Android/index.html index ea6c2c470..72f76ad53 100644 --- a/categories/Android/index.html +++ b/categories/Android/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -323,7 +323,7 @@

    Android
    - 61 + 62 日志
    diff --git a/categories/CUDA/index.html b/categories/CUDA/index.html index cb0e12818..04e469162 100644 --- a/categories/CUDA/index.html +++ b/categories/CUDA/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -340,7 +340,7 @@

    CUDA
    - 61 + 62 日志
    diff --git a/categories/Database/index.html b/categories/Database/index.html index fd1f9f1da..6c27bd21b 100644 --- a/categories/Database/index.html +++ b/categories/Database/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -346,7 +346,7 @@

    Database
    - 61 + 62 日志
    diff --git a/categories/DeepLearning/index.html b/categories/DeepLearning/index.html index 8e7991b0d..d5a917148 100644 --- a/categories/DeepLearning/index.html +++ b/categories/DeepLearning/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -207,6 +207,26 @@

    DeepLearning 2023

    + +
    @@ -363,7 +383,7 @@

    DeepLearning
    - 61 + 62 日志
    diff --git a/categories/Environment/index.html b/categories/Environment/index.html index 977902ed6..d582b3ba7 100644 --- a/categories/Environment/index.html +++ b/categories/Environment/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -363,7 +363,7 @@

    Environment
    - 61 + 62 日志
    diff --git a/categories/Life/index.html b/categories/Life/index.html index 38b062d8b..c83d7bcc7 100644 --- a/categories/Life/index.html +++ b/categories/Life/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -343,7 +343,7 @@

    Life
    - 61 + 62 日志
    diff --git a/categories/Linux/index.html b/categories/Linux/index.html index 84bb39132..32e041915 100644 --- a/categories/Linux/index.html +++ b/categories/Linux/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -343,7 +343,7 @@

    Linux
    - 61 + 62 日志
    diff --git a/categories/MachineLearning/index.html b/categories/MachineLearning/index.html index 75baafcbf..b97e84b31 100644 --- a/categories/MachineLearning/index.html +++ b/categories/MachineLearning/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -300,7 +300,7 @@

    MachineLearning
    - 61 + 62 日志
    diff --git a/categories/NLP/index.html b/categories/NLP/index.html index 2d9a2ac5c..71ffccdc8 100644 --- a/categories/NLP/index.html +++ b/categories/NLP/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -360,7 +360,7 @@

    NLP
    - 61 + 62 日志
    diff --git a/categories/Note/index.html b/categories/Note/index.html index 49e7db1a7..3033f7001 100644 --- a/categories/Note/index.html +++ b/categories/Note/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -492,7 +492,7 @@

    Note
    - 61 + 62 日志
    diff --git a/categories/PaperReading/index.html b/categories/PaperReading/index.html index ab5196040..dc8e6c921 100644 --- a/categories/PaperReading/index.html +++ b/categories/PaperReading/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -320,7 +320,7 @@

    PaperReading
    - 61 + 62 日志
    diff --git a/categories/Python/index.html b/categories/Python/index.html index 2097f02db..41990bcc9 100644 --- a/categories/Python/index.html +++ b/categories/Python/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -285,8 +285,8 @@

    Python

    -
    @@ -305,8 +305,8 @@

    Python

    -
    @@ -406,7 +406,7 @@

    Python
    - 61 + 62 日志
    diff --git a/categories/Raspberry/index.html b/categories/Raspberry/index.html index 6a399c596..39d521b7e 100644 --- a/categories/Raspberry/index.html +++ b/categories/Raspberry/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -320,7 +320,7 @@

    Raspberry
    - 61 + 62 日志
    diff --git a/categories/Scrapy/index.html b/categories/Scrapy/index.html index 3a7ad6d25..7922b34bb 100644 --- a/categories/Scrapy/index.html +++ b/categories/Scrapy/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -300,7 +300,7 @@

    Scrapy
    - 61 + 62 日志
    diff --git a/categories/Software/index.html b/categories/Software/index.html index 82a95deb2..122713214 100644 --- a/categories/Software/index.html +++ b/categories/Software/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -389,7 +389,7 @@

    Software
    - 61 + 62 日志
    diff --git a/categories/Visualization/index.html b/categories/Visualization/index.html index c77e73189..091cdaff3 100644 --- a/categories/Visualization/index.html +++ b/categories/Visualization/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -323,7 +323,7 @@

    Visualization
    - 61 + 62 日志
    diff --git a/categories/index.html b/categories/index.html index 9d26f28ca..c6e68f9a6 100644 --- a/categories/index.html +++ b/categories/index.html @@ -88,7 +88,7 @@
    -
    +
    @@ -219,7 +219,7 @@

    分类 目前共计 18 个分类

    - +
    @@ -298,7 +298,7 @@

    分类
    - 61 + 62 日志
    diff --git a/categories/test/index.html b/categories/test/index.html index 8ae9fb99d..5aed11ff1 100644 --- a/categories/test/index.html +++ b/categories/test/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -297,7 +297,7 @@

    test
    - 61 + 62 日志
    diff --git a/css/main.css b/css/main.css index 9b8e9c791..28ac555a6 100644 --- a/css/main.css +++ b/css/main.css @@ -1165,7 +1165,7 @@ pre .javascript .function { } .links-of-author a::before, .links-of-author span.exturl::before { - background: #82eb79; + background: #75c6a2; border-radius: 50%; content: ' '; display: inline-block; diff --git a/index.html b/index.html index b818e645c..24040dbdd 100644 --- a/index.html +++ b/index.html @@ -87,7 +87,7 @@
    -
    +
    @@ -194,7 +194,7 @@

    Wnma's Blogs

    - +