Merge branch 'develop' of https://github.com/PaddleJitLab/CUDATutorial into docs

Author: AndSonder

.pre-commit-config.yaml

@@ -0,0 +1,9 @@
+repos:
+  - repo: local
+    hooks:
+      - id: lint-md
+        name: lint-md
+        entry: lint-md -f
+        language: node
+        files: \.md$
+        types: [text]

docs/00_prev_concept/README.md

@@ -74,7 +74,7 @@ Build cuda_11.2.r11.2/compiler.29558016_0
 ```
 
 如果还不是很清楚 CUDA Toolkit 是什么，可以翻阅 [Nivida 官网的介绍](https://developer.nvidia.com/cuda-toolkit)：
-```
+```plain
 The NVIDIA® CUDA® Toolkit provides a development environment for creating high performance GPU-accelerated applications. With the CUDA Toolkit, you can develop, optimize, and deploy your applications on GPU-accelerated embedded systems, desktop workstations, enterprise data centers, cloud-based platforms and HPC supercomputers. The toolkit includes GPU-accelerated libraries, debugging and optimization tools, a C/C++ compiler, and a runtime library to deploy your application.
 ```
 
@@ -117,11 +117,11 @@ Hello world, CUDA! 0
 
 Windows 环境配置同样需要安装 CUDA Toolkit，下载地址为：https://developer.nvidia.com/cuda-downloads。
 
-在Windows进行安装时需要选自定义模式，采用精简模式安装后无法运行nvcc命令。
+在 Windows 进行安装时需要选自定义模式，采用精简模式安装后无法运行 nvcc 命令。
 
-安装成功后可尝试 `nvcc --version` 检测下，编译时如果缺少cl.exe，则需要安装Microsoft Visual Studio(使用C++的桌面开发版本即可)。安装完成后，将cl.exe所在路径添加到系统变量中，cl.exe通常所在文件夹为`C:\Program Files (x86)\Microsoft Visual Studio\2019\BuildTools\VC\Tools\MSVC\{version}\bin\Hostx64\x64`，具体路径根据实际安装情况有所不同。
+安装成功后可尝试 `nvcc --version` 检测下，编译时如果缺少 cl.exe，则需要安装 Microsoft Visual Studio(使用 C++的桌面开发版本即可)。安装完成后，将 cl.exe 所在路径添加到系统变量中，cl.exe 通常所在文件夹为`C:\Program Files (x86)\Microsoft Visual Studio\2019\BuildTools\VC\Tools\MSVC\{version}\bin\Hostx64\x64`，具体路径根据实际安装情况有所不同。
 
-和 Linux 不同之处在于，安装 Toolkit 之后还需要配置下环境变量。默认系统会已经有 `CUDA_PATH` 和 `CUDA_PATH_V11.0`（11.0应该是版本号），需要自己在额外添加如下环境变量：
+和 Linux 不同之处在于，安装 Toolkit 之后还需要配置下环境变量。默认系统会已经有 `CUDA_PATH` 和 `CUDA_PATH_V11.0`（11.0 应该是版本号），需要自己在额外添加如下环境变量：
 
 ```bash
 CUDA_BIN_PATH: %CUDA_PATH%\bin

docs/01_build_dev_env/README.md

@@ -51,9 +51,9 @@ GPU activities:   92.23%  570.25ms        1  570.25ms  570.25ms  570.25ms  add_k
 
 可以看出我们写的 CUDA 程序在 GPU 上主要包括 3 个关键活动：
 
-+ `add_kernel`：即执行 kernel 的时间，占比92%，耗时 570.25 ms
++ `add_kernel`：即执行 kernel 的时间，占比 92%，耗时 570.25 ms
 + HtoD 的内存拷贝：即输入 `x` → `cuda_x`，`y` → `cuda_y` 的 2 次拷贝，占比 2.99%，耗时 18.459 ms
-+ DtoH 的内存拷贝：即输出 `cuda_out` → `out` 的 1次拷贝，占比 4.79%，耗时 29.586ms 
++ DtoH 的内存拷贝：即输出 `cuda_out` → `out` 的 1 次拷贝，占比 4.79%，耗时 29.586ms 
 
 
 第三个部分是 CUDA API 的具体调用开销，这个是从 API 层面来解读各个阶段的耗时：
@@ -74,7 +74,7 @@ GPU activities:   92.23%  570.25ms        1  570.25ms  570.25ms  570.25ms  add_k
                     0.00%     543ns         1     543ns     543ns     543ns  cuDeviceGetUuid
 ```
 
-其中最耗时的就是 3 次 `cudaMemcpy` 和`cudasMalloc` 的调用，99% 的时间都在干这两个事情，可以看出显存分配是一个比较「重」的操作，任何时候我们都应该尽量避免频繁的显存分配操作。在深度学习框架中，常会借助「内存池」技术一次申请较大的显存块，然后自己管理切分、分配和回收，这样就可以减少向系统 `cudaMalloc` 的次数，感兴趣的同学可以参考[Paddle源码之内存管理技术](https://www.cnblogs.com/CocoML/p/14105729.html)。
+其中最耗时的就是 3 次 `cudaMemcpy` 和`cudasMalloc` 的调用，99% 的时间都在干这两个事情，可以看出显存分配是一个比较「重」的操作，任何时候我们都应该尽量避免频繁的显存分配操作。在深度学习框架中，常会借助「内存池」技术一次申请较大的显存块，然后自己管理切分、分配和回收，这样就可以减少向系统 `cudaMalloc` 的次数，感兴趣的同学可以参考[Paddle 源码之内存管理技术](https://www.cnblogs.com/CocoML/p/14105729.html)。
 
 剩下的 API 调用的开销基本差别不是特别大，大多数都是在 us 级别，我们一一介绍各个 API 的作用：
 

docs/03_nvprof_usage/README.md

@@ -18,7 +18,7 @@ int main(){
 }
 ```
 
-从代码中可以看出，我们是通过`<<<1, 1>>>`方式拉起了 `add_kernel` ，意味着我们只是利用了1个线程，从起点 0 &rarr; n 依次遍历张量数据做加法操作，这个完全没有利用 GPU 并行计算的优势。
+从代码中可以看出，我们是通过`<<<1, 1>>>`方式拉起了 `add_kernel` ，意味着我们只是利用了 1 个线程，从起点 0 &rarr; n 依次遍历张量数据做加法操作，这个完全没有利用 GPU 并行计算的优势。
 
 这里我们再复习下 `<<<M, T>>>` 的含义：三尖括号告诉 CUDA 在使用多少个 thread 拉起 kernel。多个线程一组成为 `thread block`，多个`thread block`一组成为 `grid`。因为前面的 `M` 表示一个 `grid` 有 `M` 个 `thread block`， 一个 `thread block` 里有 `T` 个 `thread`。
 
@@ -31,7 +31,7 @@ int main(){
 
 对于 `add_kernel<<<1, 256>>>` 而言，`threadIdx.x` 取值为 0~256 中的某个值，`blockDim.x` 的值为 256。
 
-如果要将 `N = 10000000` 切分到256个线程里并行去计算，需要调整下 `add_kernel` 中的 for 语句的写法，实现同一份代码在被不同线程调用时，自动地各自计算各自的数据，首先改成如下的范式：
+如果要将 `N = 10000000` 切分到 256 个线程里并行去计算，需要调整下 `add_kernel` 中的 for 语句的写法，实现同一份代码在被不同线程调用时，自动地各自计算各自的数据，首先改成如下的范式：
 ```cpp
 __global__ void add_kernel(float *x, float *y, float *out, int n){
     int index = 0;
@@ -62,7 +62,7 @@ __global__ void add_kernel(float *x, float *y, float *out, int n){
 完整的代码可以参考文件：[vector_add_thread.cu](./vector_add_thread.cu)。
 
 编译命令：`nvcc ./vector_add_thread.cu -o add_thread`，执行和 Profile 命令： `nvprof ./add_thread`，结果如下：
-```
+```plain
 ==36546== Profiling application: ./add_p
 ==36546== Profiling result:
             Type  Time(%)      Time     Calls       Avg       Min       Max  Name
@@ -81,7 +81,7 @@ __global__ void add_kernel(float *x, float *y, float *out, int n){
 
 ## 2. 多网格计算
 
-上面我们只用到了1个 Thread block 就实现了 38 倍的加速比。接下来我们再看如何改为多个 Thread block 的版本。
+上面我们只用到了 1 个 Thread block 就实现了 38 倍的加速比。接下来我们再看如何改为多个 Thread block 的版本。
 
 一般而言，GPU 显卡上包含了很多流式处理器（即 Streaming Multiprocessors，简称 SMs），其中每个 SM 都包含了多个并行处理单元，均支持并发地执行多个 thread block。只有将 Kernel 放在多个 thread block 上去执行，才能最大限度地利用 GPU 并行加速的能力。
 
@@ -91,7 +91,7 @@ __global__ void add_kernel(float *x, float *y, float *out, int n){
 + `gridDim.x`: 指网格（grid）的大小（size）
 
 
-同样的，我们么只需要修改 `add_kernel` 中的实现，确保每个 thread 都各自独立做计算即可，我们期望每个线程都只做1个浮点数的加法操作，也就意味着我们期望在 N = 10000000 个线程上并发地同时计算，切分示意图如下：
+同样的，我们么只需要修改 `add_kernel` 中的实现，确保每个 thread 都各自独立做计算即可，我们期望每个线程都只做 1 个浮点数的加法操作，也就意味着我们期望在 N = 10000000 个线程上并发地同时计算，切分示意图如下：
 
 ![多网格](./img/parallel_block.png)
 
@@ -115,7 +115,7 @@ __global__ void add_kernel(float *x, float *y, float *out, int n){
 完整的代码可以参考文件：[vector_add_grid.cu](./vector_add_grid.cu)。
 
 编译命令：`nvcc ./vector_add_grid.cu -o add_grid`，执行和 Profile 命令： `nvprof ./add_grid`，结果如下：
-```
+```plain
 ==32660== Profiling application: ./add_g
 ==32660== Profiling result:
             Type  Time(%)      Time     Calls       Avg       Min       Max  Name
@@ -130,7 +130,7 @@ __global__ void add_kernel(float *x, float *y, float *out, int n){
 |:---:|:---:|:---:|
 |单线程| 570 ms | - |
 |多线程| 14.7 ms| 38.7x |
-|多block| 0.153 ms| 3725x |
+|多 block| 0.153 ms| 3725x |
 
 
 ## 附参考文档

docs/04_first_refine_kernel/README.md

@@ -4,10 +4,10 @@
 
 这里我们系统性地学习下 CUDA 编程中的 Thread、Block、Grid 的概念。
 
-GPU 上一般包含很多流式处理器 SM，每个 SM 是 CUDA 架构中的基本计算单元，其可分为若干（如2~3）个网格，每个网格内包含若干（如65535）个线程块，每个线程块包含若干（如512）个线程，概要地理解的话：
+GPU 上一般包含很多流式处理器 SM，每个 SM 是 CUDA 架构中的基本计算单元，其可分为若干（如 2~3）个网格，每个网格内包含若干（如 65535）个线程块，每个线程块包含若干（如 512）个线程，概要地理解的话：
 
 + `Thread`: 一个 CUDA Kernel 可以被多个 threads 来执行
-+ `Block`: 多个 threads 会组成一个 Block，而同一个 block 中的 threads 可以同步，也可以通过shared memory通信
++ `Block`: 多个 threads 会组成一个 Block，而同一个 block 中的 threads 可以同步，也可以通过 shared memory 通信
 + `Grid`: 多个 blocks 可以组成一个 Grid
 
 其中，一个 Grid 可以包含多个 Blocks。Blocks 的分布方式可以是一维的，二维，三维的；Block 包含多个 Threads，Threads 的分布方式也可以是一维，二维，三维的。

docs/05_intro_parallel/README.md

@@ -4,15 +4,15 @@
 
 ## CUDA 层次结构
 
-在CUDA编程模型中，计算按照三级层次进行排序。每次调用CUDA内核都会创建一个新的网格，该网格由多个块组成。每个块由最多1024个单独的线程组成。这些常数可以在CUDA编程指南中查找。处于同一块内的线程可以访问相同的共享内存区域（SMEM）。
+在 CUDA 编程模型中，计算按照三级层次进行排序。每次调用 CUDA 内核都会创建一个新的网格，该网格由多个块组成。每个块由最多 1024 个单独的线程组成。这些常数可以在 CUDA 编程指南中查找。处于同一块内的线程可以访问相同的共享内存区域（SMEM）。
 
 块中的线程数可以使用一个通常称为 `blockDim` 的变量进行配置，它是一个由三个整数组成的向量。该向量的条目指定了 `blockDim.x`、`blockDim.y` 和 `blockDim.z` 的大小，如下图所示：
 
 ![picture 0](images/0b35adb64a964e56018dc9fb7277269a3efa72b1526058609e0860f33e00426b.png)  
 
 同样，网格中的块数可以使用 `gridDim` 变量进行配置。当我们从主机启动一个新的内核时，它会创建一个包含按照指定方式排列的块和线程的单一网格。
 
-对于我们的第一个内核，我们将使用 `grid`、`block` 和 `threa` 的层次结构，每个线程计算结果矩阵C中的一个元素。该线程将计算矩阵A相应行和矩阵B相应列的点积，并将结果写入矩阵C。由于矩阵C的每个位置仅由一个线程写入，我们无需进行同步。我们将以以下方式启动内核：
+对于我们的第一个内核，我们将使用 `grid`、`block` 和 `threa` 的层次结构，每个线程计算结果矩阵 C 中的一个元素。该线程将计算矩阵 A 相应行和矩阵 B 相应列的点积，并将结果写入矩阵 C。由于矩阵 C 的每个位置仅由一个线程写入，我们无需进行同步。我们将以以下方式启动内核：
 
 ```cpp
 #define CEIL_DIV(M, N) (((M) + (N)-1) / (N))
@@ -25,7 +25,7 @@ sgemm_naive<<<gridDim, blockDim>>>(M, N, K, alpha, A, B, beta, C);
 
 ## 内核实现
 
-CUDA代码是从单线程的视角编写的。在内核代码中，我们使用 `blockIdx` 和 `threadIdx`。这些变量的值会根据访问它们的线程而异。在我们的例子中，`threadIdx.x` 和 `threadIdx.y` 将根据线程在网格中的位置从0到31变化。同样，`blockIdx.x` 和 `blockIdx.y` 也将根据线程块在网格中的位置从0到 `CEIL_DIV(N, 32)` 或 `CEIL_DIV(M, 32)` 变化。
+CUDA 代码是从单线程的视角编写的。在内核代码中，我们使用 `blockIdx` 和 `threadIdx`。这些变量的值会根据访问它们的线程而异。在我们的例子中，`threadIdx.x` 和 `threadIdx.y` 将根据线程在网格中的位置从 0 到 31 变化。同样，`blockIdx.x` 和 `blockIdx.y` 也将根据线程块在网格中的位置从 0 到 `CEIL_DIV(N, 32)` 或 `CEIL_DIV(M, 32)` 变化。
 
 ```cpp
 __global__ void sgemm_naive_kernel(float *A, float *B, float *C, int M, int N, int K)
@@ -50,7 +50,7 @@ __global__ void sgemm_naive_kernel(float *A, float *B, float *C, int M, int N, i
 
 运行命令：
 
-```
+```plain
 nvcc -o matmul_raw matmul_raw.cu
 ./matmul_raw
 ```