fix(accelerator/practice): fix typo/error (#364)

* fix(accelerator/practise): fix typo & ref

* fix(accelerator/practise): fix duplicated_data.svg

* fix(accelerator/practise):fix text

Co-authored-by: wenteng_liang@163.com <Went-Liang>

chapter_accelerator/accelerator_practise.md

@@ -217,7 +217,7 @@ Average Throughput: 185.313 GFLOPS
 
 #### 使用封装结构代替指针
 
-在上面的实现中，由于二维矩阵的数据是使用一维数组进行存储，所以在访问数据时需要使用行坐标与二维矩阵宽度的乘积和列坐标的和来索引到具体位置的元素，这样的访问方式并不直观且在后续逐渐复杂的实现中容易出错。因此，我们可以自定义一个结构体，通过重载 `()` 运算符，实现对矩阵元素的二维索引。而且在后续的实现中，在访问数据块内的数据时，我们可能使用大量重复的坐标索引偏移，为了简化代码，我们提供 `addOffset` 方法用于加入一个固定的偏移，具体实现如下：
+在上面的实现中，由于二维矩阵的数据是使用一维数组进行存储，所以在访问数据时需要使用行坐标与二维矩阵宽度的乘积和列坐标的和来索引到具体位置的元素，这样的访问方式并不直观且在后续逐渐复杂的实现中容易出错。因此，我们可以自定义一个结构体，通过重载 `()` 运算符，实现对矩阵元素的二维索引，同时我们提供 `addOffset` 方法用于加入一个固定的偏移，具体实现如下：
 
 ```c++
 template <typename T>
@@ -291,7 +291,7 @@ __global__ void gemmKernel(const float * A,
 
 计算强度（Compute Intensity）指计算指令数量与访存指令数量的比值，在现代GPU中往往有大量计算单元但只有有限的访存带宽，程序很容易出现计算单元等待数据读取的问题，因此提高计算强度是提升程序性能的一条切实有限的指导思路。对于之前实现的GPU核函数，我们可以粗略计算其计算强度：在$K$次循环的内积计算中，对矩阵$A$与矩阵$B$的每次读取会计算一次浮点乘法与浮点加法，因此计算强度为1——两次浮点运算除以两次数据读取。之前的版本是每个线程负责处理矩阵$C$的一个元素——计算矩阵$A$的一行与矩阵$B$的一列的内积，我们可以通过使每个线程计算$C$更多的元素——计算矩阵$A$的多行与矩阵$B$的多列的内积——从而提升计算强度。具体地，如果在$K$次循环的内积计算中一次读取矩阵$A$中的$m$个元素和矩阵$B$中的$n$个元素，那么访存指令为$m+n$条，而计算指令为$2mn$条，所以计算强度为$\frac{2mn}{m+n}$，因此可以很容易发现提高$m$和$n$会带来计算强度的提升。
 
-我们在上一个代码例子中对全局内存的访问与存储都是借助 `float` 指针完成的，具体到硬件指令集上实际是使用指令 `LDG.E` 与 `STG.E` 完成的。我们可以使用128位宽指令`LDG.E.128` 与 `STG.E.128`  一次读取多个 `float` 数。使用宽指令的好处一方面有简化了指令序列，使用一个宽指令代替四个标准指令可以节省十几个指令的发射周期，这可以为计算指令的发射争取到额外的时间；此外128比特正好等于一个cache line的长度，使用宽指令也有助于提高cache line的命中率。但我们并不提倡在一切代码中过度追求宽指令的使用，开发者应当将更多的时间关注并行性设计和局部数据复用等更直接的优化手段。
+我们在上一个代码例子中对全局内存的访问与存储都是借助 `float` 指针完成的，具体到硬件指令集上实际是使用指令 `LDG.E` 与 `STG.E` 完成的。我们可以使用128位宽指令`LDG.E.128` 与 `STG.E.128`  一次读取多个 `float` 数。使用宽指令的好处是一方面简化了指令序列，使用一个宽指令代替四个标准指令可以节省十几个指令的发射周期，这可以为计算指令的发射争取到额外的时间；另一方面128比特正好等于一个cache line的长度，使用宽指令也有助于提高cache line的命中率。但我们并不提倡在一切代码中过度追求宽指令的使用，开发者应当将更多的时间关注并行性设计和局部数据复用等更直接的优化手段。
 
 具体的实现如下，由于每个 `float` 类型大小为32个比特，我们可以将4个 `float` 堆叠在一起构成一个128比特的 `float4` 类，对 `float4` 的访存将会是使用宽指令完成。虽然CUDA Toolkit已经有实现的 `float4` 类，但是为了代码抽象我们将自行实现我们自己的 `float4` 类。
 
@@ -363,7 +363,7 @@ __global__ void gemmKernel(const float *__restrict__ A,
 }
 ```
 
-我们首先在第6到14行计算每个线程需要处理的数据块在矩阵中的起始行列坐标`m,n`，即图2中矩阵$C$中绿色数据块的左上角坐标，然后使用`Tensor2D`中的`addOffset`方法，为每个线程定位到它要处理的数据块的起始位置上，并且利用`validOffset`方法判断线程是否越界。然后就可以沿着K方向循环，在第18到23行每个线程分别读取矩阵$A$中连续的4行和矩阵$B$中连续的四列组成两个 `float4` ，即图2中粉色与黄色的4维向量。之后在第25到27行计算线程负责处理矩阵$C$中的$4 \times 4$个元素。最后在第30到40行对结果使用参数 `alpha` 和 `beta` 进行放缩并写回矩阵$C$的内存。
+我们首先在第6到14行计算每个线程需要处理的数据块在矩阵中的起始行列坐标`m,n`，即 :numref:`use_float4` 中矩阵$C$中浅绿色数据块的左上角坐标，然后使用`Tensor2D`中的`addOffset`方法，为每个线程定位到它要处理的数据块的起始位置上，并且利用`validOffset`方法判断线程是否越界。然后就可以沿着K方向循环，在第18到23行每个线程分别读取矩阵$A$中连续的4行和矩阵$B$中连续的四列组成两个 `float4` ，即 :numref:`use_float4` 中粉色与黄色的四个元素。之后在第25到27行计算线程负责处理矩阵$C$中的$4 \times 4$个元素。最后在第30到40行对结果使用参数 `alpha` 和 `beta` 进行放缩并写回矩阵$C$的内存。
 
 ![提高计算强度](../img/ch06/6.4/use_float4.svg)
 :width:` 800px`
@@ -430,7 +430,7 @@ ncu --set full -o first_attepmt_prof_result ./first_attempt
 
 ### 进一步提升计算强度
 
-我们可以通过使每个线程负责处理更多的矩阵$C$中的数据块从而实现更高的计算强度，即如下图3右侧所示，使 `thread tile` 扩大为4个$4 \times 4$矩阵的规模。我们对核函数进行以下修改，首先我们用`LayoutTile` 来描述每个线程块处理数据 `tile`的布局 ，其中 `LayoutTile::m` 和 `LayoutTile::n` 等于 :numref:`use_tile` 左图中浅绿色矩阵块的高度和宽度， `LayoutTile::k` 等于1；其次我们用`LayoutBlock` 来描述一个线程块中线程的布局；同时我们用`LayoutThread` 来描述 `thread tile` 中子矩阵的布局 ，其中`LayoutThread::m` 和 `LayoutThread::n` 等于 :numref:`use_tile` 右图中深绿色矩阵块的高度和宽度 。
+我们可以通过使每个线程负责处理更多的矩阵$C$中的数据块从而实现更高的计算强度，即如 :numref:`use_tile` 右侧所示，使 `thread tile` 扩大为4个$4 \times 4$矩阵的规模。我们对核函数进行以下修改，首先我们用`LayoutTile` 来描述每个线程块处理数据 `tile`的布局 ，其中 `LayoutTile::m` 和 `LayoutTile::n` 等于 :numref:`use_tile` 左图中浅绿色矩阵块的高度和宽度， `LayoutTile::k` 等于1；其次我们用`LayoutBlock` 来描述一个线程块中线程的布局；同时我们用`LayoutThread` 来描述 `thread tile` 中子矩阵的布局 ，其中`LayoutThread::m` 和 `LayoutThread::n` 等于 :numref:`use_tile` 右图中深绿色矩阵块的高度和宽度 。
 
 ![通过提高线程所处理矩阵块的数量来进一步提高计算强度](../img/ch06/6.4/use_tile.svg)
 :width:` 800px`
@@ -464,7 +464,7 @@ const unsigned intervalA = LayoutTile::m / iterationA;
 const unsigned intervalB = LayoutTile::n / iterationB;
 ```
  `iterationA` 是每个线程处理 `thread tile` 在行方向上迭代的次数。`intervalA` 是 `thread tile` 子矩阵在行方向的间隔。同理 `iterationB` 与 `intervalB` 是在列方向上数据块的数量与数据块的间隔。
-因为 `thread tile` 扩大为若干个矩阵块，我们每个线程在读取数据前需要判断地址是否越界。 我们使用以下代码用来记录每个矩阵块是否越界：
+因为 `thread tile` 扩大为若干个矩阵块，我们使用以下代码用来记录每个矩阵块是否越界：
 
 ```c++
 bool validLoadTileA[iterationA];
@@ -506,7 +506,7 @@ for (unsigned k = 0; k < K; ++k) {
   }
 }
 ```
-注意到我们此时使用了编译器指令 `#pragma unroll` 用于将循环展开，即如果循环次数是可以在编译时确定的话，编译器将会把带有判断和跳转的循环代码展开成串行代码。这样做的好处主要是减少了判断语句，此外还可以便于编译器发现数据依赖从而更好地分配寄存器。缺点是会增加寄存器的使用，有潜在的降低GPU占用率的风险。
+注意到我们此时使用了编译器指令 `#pragma unroll` 用于将循环展开，即如果循环次数是可以在编译时确定的话，编译器将会把带有判断和跳转的循环代码展开成串行代码。这样做的好处主要是减少了判断语句，此外还有利于编译器发现数据依赖从而更好地分配寄存器。缺点是可能会增加寄存器的使用，有潜在的降低GPU占用率的风险。
 最后对于结果使用 `alpha` 和 `beta` 的放缩以及写回也相应的加上数据块的循环：
 
 ```c++
@@ -554,13 +554,13 @@ Average Time: 3.188 ms, Average Throughput: 2694.440 GFLOPS
 
 :label:`sec-accelerator-use-smem`
 
-虽然令一个线程一次读取更多的数据能取得计算强度的提升进而带来性能的提升，但是这种设计会导致由于单个线程处理数据的增多导致开启总的线程数量减少，进而导致并行度下降，因此我们需要使用其他硬件特性在尽可能不影响并行度的前提下取得性能提升。在之前的代码中，我们开启若干个线程块，每个线程块处理矩阵$C$中的一个或多个矩阵块。在 :numref:`duplicated_data` 中，我们可以观察到，处理矩阵$C$同一行的线程$x, y$会读取矩阵$A$中相同的数据，我们可以借助共享内存让同一个线程块中不同的线程读取不重复的数据而提升程序吞吐量。
+虽然令一个线程一次读取更多的数据能取得计算强度的提升进而带来性能的提升，但是这种令单个线程处理数据增多的设计会导致开启总的线程数量减少，进而导致并行度下降，因此我们需要使用其他硬件特性在尽可能不影响并行度的前提下取得性能提升。在之前的代码中，我们开启若干个线程块，每个线程块处理矩阵$C$中的一个或多个矩阵块。在 :numref:`duplicated_data` 中，我们可以观察到，处理矩阵$C$同一行的线程$x, y$会读取矩阵$A$中相同的数据，我们可以借助共享内存让同一个线程块中不同的线程读取不重复的数据而提升程序吞吐量。
 
 ![线程间重复读取数据](../img/ch06/6.4/duplicated_data.svg)
 :width:` 800px`
 :label:`duplicated_data`
 
-具体地，我们需要对代码进行如下改造：首先此前代码在计算内积过程是分$K$次循环读取数据并累加计算，在此设定下每次循环中处理矩阵$C$中相同行的线程会读取相同的矩阵$A$的数据，处理矩阵$C$中相同列的线程会读取相同的矩阵$B$的数据。我们可以通过将此$K$次循环拆解成两层循环，外层循环$\frac{K}{tileK}$次，每次外循环的迭代读取一整块数据，内层循环$tileK$次进行累加数据。直观来看，外层循环如 :numref:`use_smem_store` 所示，每次循环将矩阵$A$和矩阵$B$中一整个 `tile` 读取到共享内存中；内层循环如 :numref:`use_smem_load` 所示，每次循环从共享内存读取数据并计算。这种设计带来的好处是，我们可以让每个线程不必独自从全局内存读取所有需要的数据，整个线程块将共同需要的数据从全局内存中读取并写入到共享内存中，此后每个线程在计算过程中只需要从共享内存中读取所需要的数据即可。
+具体地，我们需要对代码进行如下改造：首先此前代码在计算内积过程是进行$K$次循环读取数据并累加计算，在此设定下每次循环中处理矩阵$C$中相同行的线程会读取相同的矩阵$A$的数据，处理矩阵$C$中相同列的线程会读取相同的矩阵$B$的数据。我们可以通过将此$K$次循环拆解成两层循环，外层循环$\frac{K}{tileK}$次，每次外循环的迭代读取一整块数据，内层循环$tileK$次进行累加数据。直观来看，外层循环如 :numref:`use_smem_store` 所示，每次循环将矩阵$A$和矩阵$B$中一整个 `tile` 读取到共享内存中；内层循环如 :numref:`use_smem_load` 所示，每次循环从共享内存读取数据并计算。这种设计带来的好处是，我们可以让每个线程不必独自从全局内存读取所有需要的数据，整个线程块将共同需要的数据从全局内存中读取并写入到共享内存中，此后每个线程在计算过程中只需要从共享内存中读取所需要的数据即可。
 
 
 ![向共享内存中写入数据](../img/ch06/6.4/use_smem_store.svg)
@@ -571,7 +571,7 @@ Average Time: 3.188 ms, Average Throughput: 2694.440 GFLOPS
 :width:` 800px`
 :label:`use_smem_load`
 
-下面我们将实现使用共享内存的GPU核函数。首先，我们可以计算每个线程块在外层循环的每次迭代中从矩阵$A$中读取大小为$tileM \times tileK$的数据块，在矩阵$B$中读取大小为$tileK \times tileN$的数据块。假设每个线程块中一共含有$blockSize$个线程，那么就可以使用这$blockSize$个线程，每个线程循环$\frac{tileM * tileK}{blockSize * 4}$次将矩阵$A$中的矩阵块 `tileA` 读取进共享内存中，同理每个线程循环$\frac{tileM * tileK}{blockSize * 4}$次将矩阵$B$中的矩阵块 `tileB` 读取进共享内存中。
+下面我们将实现使用共享内存的GPU核函数。首先，我们定义每个线程块在外层循环的每次迭代中从矩阵$A$中读取大小为$tileM \times tileK$的数据块，在矩阵$B$中读取大小为$tileK \times tileN$的数据块。假设每个线程块中一共含有$blockSize$个线程，那么就可以使用这$blockSize$个线程，每个线程循环$\frac{tileM * tileK}{blockSize * 4}$次将矩阵$A$中的矩阵块 `tileA` 读取进共享内存中，同理每个线程循环$\frac{tileM * tileK}{blockSize * 4}$次将矩阵$B$中的矩阵块 `tileB` 读取进共享内存中。
 
 首先需要定义若干变量：
 
@@ -599,7 +599,8 @@ constexpr unsigned tileSizeB = LayoutTile::n * LayoutTile::k;
 constexpr unsigned tileIterationsB = tileSizeB / blockSize / ratio;
 constexpr unsigned tileGlobalIntervalB = blockSize / LayoutTileT::n;
 constexpr unsigned tileComputeIterationsB = LayoutTileT::n / LayoutBlock::n;
-constexpr unsigned tileSharedIntervalBT = LayoutTileT::n / tileComputeIterationsB;const unsigned nInTileB = threadIdx.x % LayoutTileT::n;
+constexpr unsigned tileSharedIntervalBT = LayoutTileT::n / tileComputeIterationsB;
+const unsigned nInTileB = threadIdx.x % LayoutTileT::n;
 const unsigned kinTileB = threadIdx.x / LayoutTileT::n;
 ```
 因为 `LayoutTile` 与 `LayoutThread` 是表示的 `float` 数据的布局，我们有时将其看为 `float4` 的数据储存，因此我们需要加入变量 `LayoutTileT` 与 `LayoutThreadT` 。 `blockSize` 指一个线程块内的线程数量。 我们在此版本使用一维线程块的布局模拟二维布局，所以我们需要计算在二维布局下的坐标：用 `mInTileC` 与 `nInTileC` 表示在给定 `LayoutBlock` 布局下的二维线程坐标。由于 `tileA` 是$tileM \times timeK$的尺寸，因此我们可以确定其中数据数量`tileSizeA` ，由于一个线程块内有 `blockSize` 个线程且每个线程一次读取 `ratio` 个 `float` 数，因此整个 `tileA` 需要用 `tileIterationsA = tileSizeA / blockSize / ratio` 次读取。每个线程在最开始时负责读取的 `tileA` 的位置使用变量 `kInTileA` 和 `mInTileA` 表示。因为需要用`tileIterationsA` 次读取 `tileA` ，每次向下滑动的距离我们使用变量`tileGlobalIntervalA`表示。同时因为需要用每个线程需要处理 `thread tile`  中多个子矩阵块，其中每个线程处理 `thread tile` 时在行方向上迭代的次数 定义为`tileComputeIterationsA` 。这些子矩阵块在 `m` 方向的间隔我们用`tileSharedIntervalA` 表示。类似地，我们定义与 `tileB` 的若干变量。
@@ -767,7 +768,7 @@ Average Time: 0.610 ms, Average Throughput: 14083.116 GFLOPS
 
 ### 隐藏共享内存读取延迟
 
-在GPU中读取数据共享内存中的数据使用指令 `LDS` ，在这条指令发出后并不会等待数据读取到寄存器后再执行下一条语句，只有执行到依赖 `LDS` 指令读取的数据的指令时才会等待读取的完成。而在上一小节中，我们在内层$tileK$次循环中，每次循环迭代对共享内存发射完读取指令之后就会立即执行依赖于读取数据的数学运算，这样就会导致计算单元等待数据从共享内存的读取，如 :numref:`use_smem_pipeline` 所示。事实上，对共享内存的访问周期能多达几十个时钟周期，而计算指令的执行往往只有几个时钟周期，因此通过一定方式隐藏对共享内存的访问会取得不小的收益。我们可以重新优化流水线隐藏一定的数据读取延迟。具体地，我们可以在内层的$tileK$次循环中每次循环开始时读取发射下一次内层循环数据的读取指令。由于在执行本次运算时计算指令并不依赖于下一次循环的数据，因此计算过程不会等待之前发出的读取下一次内层循环数据的指令，具体见下 :numref:`hide_smem_latency` 。
+在GPU中使用指令 `LDS` 读取共享内存中的数据，在这条指令发出后并不会等待数据读取到寄存器后再执行下一条语句，只有执行到依赖 `LDS` 指令读取的数据的指令时才会等待读取的完成。而在上一小节中，我们在内层$tileK$次循环中，每次发射完读取共享内存的指令之后就会立即执行依赖于读取数据的数学运算，这样就会导致计算单元等待数据从共享内存的读取，如 :numref:`use_smem_pipeline` 所示。事实上，对共享内存的访问周期能多达几十个时钟周期，而计算指令的执行往往只有几个时钟周期，因此通过一定方式隐藏对共享内存的访问会取得不小的收益。我们可以重新优化流水线隐藏一定的数据读取延迟。具体地，我们可以在内层的$tileK$次循环中每次循环开始时读取发射下一次内层循环数据的读取指令。由于在执行本次运算时计算指令并不依赖于下一次循环的数据，因此计算过程不会等待之前发出的读取下一次内层循环数据的指令，具体见 :numref:`hide_smem_latency` 。
 
 ![上一个GPU核函数的流水线](../img/ch06/6.4/use_smem_pipeline.svg)
 :width:` 800px`
@@ -837,11 +838,11 @@ for (unsigned d = 0; d < tileComputeIterationsA * LayoutThread::m; ++d) {
 Max Error: 0.000092
 Average Time: 0.585 ms, Average Throughput: 14686.179 GFLOPS
 ```
-使用Nsight Compute观察发现：相比上一个GPU核函数，指标 `Stall Short Scoreboard` 减少了67%。Scoreboard的功能是由于此前提过GPU内存读写指令发出后并不会等待数据读取到寄存器后再执行下一条语句，但是会在Scoreboard设置符号并在完成读取后置回符号，等到之后有数据依赖的指令执行前会等待Scoreboard中符号的置回。这里`Stall Short Scoreboard` 的减少充分说明了内存延迟是有效的。
+使用Nsight Compute观察发现：相比上一个GPU核函数，指标 `Stall Short Scoreboard` 减少了67%。而此前提过GPU内存读写指令发出后并不会等待数据读取到寄存器后再执行下一条语句，但是会在Scoreboard设置符号并在完成读取后置回符号，等到之后有数据依赖的指令执行前会等待Scoreboard中符号的置回。所以这里`Stall Short Scoreboard` 的减少充分说明了内存延迟是有效的。
 
 ### 隐藏全局内存读取延迟
 
-上一小节中我们介绍了对共享内存读取流水线优化的方法，事实上，GPU再读取全局内存中使用的指令 `LDG` 也有与共享内存读取指令 `LDS` 类似的行为特性。因此我们类似的在$\frac{K}{tileK}$次外层循环中每次循环开始时发出下一次外层循环需要的矩阵$A$中的数据块的读取指令，而本次外循环的整个内层循环过程中不依赖下一次外循环的数据，因此本次外循环的内循环过程中不会等待对下一次外层循环需要的矩阵$A$中的数据块的读取指令完成，从而实现隐藏全局内存读取延迟的目的。具体流水线可视化见 :numref:`hide_global_latency` ：
+上一小节中我们介绍了对共享内存读取流水线优化的方法，事实上，GPU再读取全局内存中使用的指令 `LDG` 也有与共享内存读取指令 `LDS` 类似的行为特性。因此我们类似的在$\frac{K}{tileK}$次外层循环中每次循环开始时发出下一次外层循环需要的矩阵$A$中的数据块的读取指令，而本次外循环的整个内层循环过程中不依赖下一次外循环的数据，因此本次外循环的内循环过程中不会等待对下一次外层循环需要的矩阵$A$中的数据块的读取指令完成，从而实现隐藏全局内存读取延迟的目的。具体流水线可视化见 :numref:`hide_global_latency` 。
 
 ![隐藏全局内存读取延迟的流水线](../img/ch06/6.4/hide_global_latency.svg)
 :width:` 800px`
@@ -945,9 +946,9 @@ Average Time: 0.542 ms, Average Throughput: 15838.302 GFLOPS
 ```
 使用Nsight Compute分析我们观察到指标 `Stall Long Scoreboard` 减少了67%，与上一小结的 `Stall Short Scoreboard` 概念相对应，`Stall Long Scoreboard` 主要是针对全局内存的指标。该指标的显著减少充分说明我们可以在一定程度上隐藏全局内存的读取。
 
-### 与cuBlas对比
+### 与cuBLAS对比
 
-前一节中介绍了cuBlas的接口，我们可以很容易地写出以下代码使用cuBlas完成矩阵乘法：
+前一节中介绍了cuBLAS的接口，我们可以很容易地写出以下代码使用cuBLAS完成矩阵乘法：
 
 ```c++
 void cublasGemm(const float *A, const float *B, float *C, float alf, float bet, int M, int N, int K) {
@@ -972,11 +973,11 @@ Max Error: 0.000092
 Average Time: 0.613 ms, Throughput: 14002.600 GFLOPS
 ```
 使用Nsight Compute分析发现 `LDG` 和 `STS` 等指令使用较多，导致指令发射压力较大，具体体现在 `Stall Wait` 与 `Stall Dispatch Stall` 指标相比我们较差。但其他指标诸如 `Stall Long Scoreboard` 等优于我们，但总体上我们略胜一筹。
-尽管我们的代码相比cuBlas已经取得了一定的性能提升，但是需要强调的是cuBlas内部为各种不同的矩阵尺寸以及不同的设备实现了若干不同的GPU核函数，我们实现的核函数在其他尺寸或其他设备设备上性能可能无法取得此加速比。
+尽管我们的代码相比cuBLAS已经取得了一定的性能提升，但是需要强调的是cuBLAS内部为各种不同的矩阵尺寸以及不同的设备实现了若干不同的GPU核函数，我们实现的核函数在其他尺寸或其他设备设备上性能可能无法取得此加速比。
 
-1.  **并行资源映射——提高并行性**：将多层级的并行资源（`block` 、`warp` 、`thread` ）与对应需要计算/搬移的数据建立映射关系，提高程序并行性。将可并行的计算/数据搬移操作映射到并行资源上，对于GEMM实例，在朴素实现的例子中 :numref:`sec-accelerator-naive` ，我们令每个`block` 与矩阵$C$中的一个矩阵块建立映射关系，每个`thread` 与矩阵块中的一个元素建立映射关系。
+1.  **并行资源映射——提高并行性**：将多层级的并行资源（`block` 、`warp` 、`thread` ）与对应需要计算/搬移的数据建立映射关系，提高程序并行性。将可并行的计算/数据搬移操作映射到并行资源上，对于一般矩阵乘法实例，在朴素实现的例子中 :numref:`sec-accelerator-naive` ，我们令每个`block` 与矩阵$C$中的一个矩阵块建立映射关系，每个`thread` 与矩阵块中的一个元素建立映射关系。
 2.  **优化内存结构——减小访存延迟**：观察计算过程中同一个`block` 中数据复用的情况，将复用的数据被如共享内存、寄存器等高性能体系结构存储下来，以此提高吞吐量。如在 :numref:`sec-accelerator-use-smem` 中我们将矩阵$A$与矩阵$B$中会被同一个`block` 内不同`thread` 共同访问的数据缓存到共享内存中。
-3.  **优化指令执行——减小指令发射开销**：使用\#unroll功能进行循环展开来提升指令级并行，减少逻辑判断；使用向量化加载指令以提高带宽等，对于Ampere架构，最大向量化加载指令为ldg128，即128比特带宽，可以采用`float4` 类型的数据进行读取。
+3.  **优化指令执行——减小指令发射开销**：使用\#unroll功能进行循环展开来提升指令级并行，减少逻辑判断；使用向量化加载指令以提高带宽等，对于Ampere架构，最大向量化加载指令为`LDG.E.128`，可以采用`float4` 类型的数据进行读取。
 4.  **优化访存流水线——隐藏访存延迟**：在进行内存结构变化（矩阵数据搬移）时，可以优化访存流水线，在数据搬移的间隔执行计算操作以隐藏数据搬移的延迟。
 
 ### 扩展
@@ -987,7 +988,7 @@ Average Time: 0.613 ms, Throughput: 14002.600 GFLOPS
 
 PTX指令集是英伟达推出的一种CUDA指令集中间表示，其并不是直接由二进制反汇编得来的汇编语言，事实上它更像是Java语言中的ByteCode概念。PTX给用户更细粒度的编程的可能，但是要注意的是由于它并不是直接的汇编语言，所以最后被CUDA汇编器汇编得到的二进制文件可能会用其他等价语句实现。例如，在安培设备中不支持 `mma.m8n8k4` 相关功能，但用户仍然可以在代码中直接使用此PTX指令，但在汇编过程会被实现为等价的非 `mma` 指令，进而导致性能不如预期。
 
-相关的PTX指令主要有 `wmma` 和 `mma` 。 `wmma` 是相对 `mma` 高一层的指令，它可以提供若干修饰符如 `load` 用于读取数据， `store` 用于储存数据， `mma` 用于计算等操作。
+与一般矩阵乘法相关的PTX指令主要有 `wmma` 和 `mma` 。 `wmma` 是相对 `mma` 高一层的指令，它可以提供若干修饰符如 `load` 用于读取数据， `store` 用于储存数据， `mma` 用于计算等操作。
 
 修饰符的使用可以类似于以下形式 `wmma.load.a.sync.aligned.m16n16k16.global.row.f16` 。 `a` 代表目标为矩阵$A$中的数据；`sync` 代表线程束内线程在执行这条指令之前会进行同步； `aligned` 代表要求线程束内所有线程都执行相同的指令； `m16n16k16` 代表一个线程束的线程共同完成数据处理的规模； `global` 代表是从全局内存读取； `row` 代表数据是行优先的布局储存的； `f16` 代表数据类型为半精度。其他可能的修饰符可以在官方文档中查阅。 `mma` 指令相比 `wmma.mma` 功能更加丰富，使用也更加灵活，其使用可类似以下形式 `mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16` 。 `sync` 代表线程束内线程在执行这条指令之前会进行同步； `aligned` 代表要求线程束内所有线程都执行相同的指令； `m16n8k16` 代表一个线程束的线程共同完成数据处理的规模； `row.col` 代表从矩阵$A$读入的矩阵块布局是行优先，从矩阵$B$读入的数据块布局是列优先；`f16.f16.f16.f16` 分别代表累加结果的矩阵块、从矩阵$A$读入的矩阵块、从矩阵$B$读入的矩阵块、累加输入的矩阵块的数据类型都是半精度浮点数。