diff --git a/docs/09_optimize_reduce/01_interleaved_addressing/reduce_interleaved_addressing.cu b/docs/09_optimize_reduce/01_interleaved_addressing/reduce_interleaved_addressing.cu index b62f9ee..c467ec3 100644 --- a/docs/09_optimize_reduce/01_interleaved_addressing/reduce_interleaved_addressing.cu +++ b/docs/09_optimize_reduce/01_interleaved_addressing/reduce_interleaved_addressing.cu @@ -25,7 +25,7 @@ __global__ void reduce_naive_kernel(int *arr, int *out, int len) for (int s = 1; s < bdim; s *= 2) { int index = 2 * s * tid; - if ((index + s < bdim) && (bdim * bid + s < len)) + if ((index + s < bdim) && (bdim * bid + s + index < len)) { sdata[index] += sdata[index + s]; } @@ -98,4 +98,4 @@ int main() delete[] arr; delete[] out; return 0; -}⏎ +} diff --git a/docs/12_convolution/02_intro_conv_optimize/README.md b/docs/12_convolution/02_intro_conv_optimize/README.md index 262fe35..5fa6b75 100644 --- a/docs/12_convolution/02_intro_conv_optimize/README.md +++ b/docs/12_convolution/02_intro_conv_optimize/README.md @@ -6,7 +6,7 @@ ## 1. 卷积算法映射为矩阵乘法 -首先我们先来回顾一下卷积算法的定义,假设输入的特征图为 $X$,卷积核为 $K$,输出特征图为 $Y$,$X$ 的大小为 $N \times C \times H \times W$,$K$ 的大小为 $M \times C \times K_h \times K_w$,$Y$ 的大小为 $N \times M \times H \times W$。那么卷积算法的定义如下: +首先我们先来回顾一下卷积算法的定义,假设输入的特征图为 $X$,卷积核为 $K$,输出特征图为 $Y$,$X$ 的大小为 $N \times C_{in} \times H_{in} \times W_{in}$,$K$ 的大小为 $M \times C_{in} \times K_h \times K_w$,$Y$ 的大小为 $N \times M \times H_{out} \times W_{out}$。那么卷积算法的定义如下: $$ Y[n,oc,oh,ow] = \sum_{ic}\sum_{fh}\sum_{fw}X[n,ic,ih,iw] \times K[oc,ic,fh,fw] diff --git a/docs/17_flash_attn/01_flash_attn_v1_part1.md b/docs/17_flash_attn/01_flash_attn_v1_part1.md index 62940ef..69bd8c2 100644 --- a/docs/17_flash_attn/01_flash_attn_v1_part1.md +++ b/docs/17_flash_attn/01_flash_attn_v1_part1.md @@ -88,7 +88,7 @@ $$ --- -**步骤 1:外层循环 $ j=1 $,处理块 $\mathbf{K}_1, \mathbf{V}_1$** +**步骤 1:外层循环 $ j=1 $,处理块 $\mathbf{K}_1, \mathbf{V}_1$ ** 1. **加载 $\mathbf{K}_1, \mathbf{V}_1$ 到 SRAM**: @@ -96,7 +96,7 @@ $$ \mathbf{K}_1 = \begin{bmatrix} k_{11} & k_{12} \\ k_{21} & k_{22} \end{bmatrix}, \quad \mathbf{V}_1 = \begin{bmatrix} v_{11} & v_{12} \\ v_{21} & v_{22} \end{bmatrix} $$ -2. **内层循环 $ i=1 $,处理块 $\mathbf{Q}_1$**: +2. **内层循环 $ i=1 $,处理块 $\mathbf{Q}_1$ **: - **加载数据**: $$ \mathbf{Q}_1 = \begin{bmatrix} q_{11} & q_{12} \\ q_{21} & q_{22} \end{bmatrix}, \quad \mathbf{O}_1 = \begin{bmatrix} 0 & 0 \\ 0 & 0 \end{bmatrix}, \quad \ell_1 = [0, 0]^T, \quad m_1 = [-\infty, -\infty]^T @@ -122,7 +122,7 @@ $$ - 类似地,加载 $\mathbf{Q}_2 = \begin{bmatrix} q_{31} & q_{32} \\ q_{41} & q_{42} \end{bmatrix}$,计算 $\mathbf{S}_{21} = \mathbf{Q}_2 \mathbf{K}_1^T$,更新后两行 $\mathbf{O}_2$。 -**步骤 2:外层循环 $ j=2 $,处理块 $\mathbf{K}_2, \mathbf{V}_2$** +**步骤 2:外层循环 $ j=2 $,处理块 $\mathbf{K}_2, \mathbf{V}_2$ ** 1. **加载 $\mathbf{K}_2, \mathbf{V}_2$ 到 SRAM**: $$ @@ -142,7 +142,7 @@ $$ $$ - **结果等价于全局 Softmax**:最终 $\mathbf{O}_1$ 为前两行注意力结果的加权和。 -3. **内层循环 $ i=2 $,处理块 $\mathbf{Q}_2$**: +3. **内层循环 $ i=2 $,处理块 $\mathbf{Q}_2$ **: - 类似地,计算 $\mathbf{S}_{22} = \mathbf{Q}_2 \mathbf{K}_2^T$,更新后两行 $\mathbf{O}_2$。