抱歉，您的浏览器无法访问本站

cutlass cute compose语义级理解layout是一个映射, 可以将逻辑(m, n)映射到物理的index: layout(m, n) -> idx。当两个layout复合时, 该idx除了有”物理idx”的还有外, 还有一层”逻辑crd”的语义。如A o B = A(B(i)) i传入到B, 先转变为B视角的逻辑crd crd经过B映射后得到一个新的idx’ 该idx...

welford算法求方差 cc https://zhuanlan.zhihu.com/p/408474710 基础方法 $D(x) = \frac{\sum(x_i - mean)^2}{n}$ 需要两次遍历1. 第一次遍历获取mean, 2. 第二次遍历计算方差 数学变换优化 数学等价转换的一次遍历 可以推导出$D(x) = E(x^2) - E(x)^2$, 只需要用x和x方做一...

Cutlass Tiled Copy Copy is all you need. make_tiled_copy语义理解。核心在于: tiler和layout_tv。先说结论: 用atom去对tv layout进行分tile。用tiler去对目标tensor进行分tile。最后将这两个layout组合得到新的tv layout，表示tile-wise的访问tv, v的layout能够保证满...

CUDA占用优化 二次吸收: https://medium.com/@manisharadwad/unlocking-gpu-potential-understanding-and-optimizing-cuda-occupancy-2f43ee01ad7e 优化问题, 由于分配的粒度问题(一个block一个block分配资源, 固定会有n_thread * reg, n_thread的整...

GPU L2 Cache优化In short: 1. 总体需要经手L2 Cache的数据少了，命中率就高了 2. 访问相同数据的thread block越多，命中率就高了 https://triton-lang.org/main/getting-started/tutorials/03-matrix-multiplication.html#l2-cache-optimizations L2...

Flash Attention v3技术点 hopper特性co-design hopper新特性 低精度 优化量化误差 异步 intra-warpgroup, inter-warpgroup 异步 异步: 异步计算(WGMMA), 异步传输(TMA) => 软件级流水线 producer-consumer模型: 生产者传输, 消费者计算 算力高的tensor-co...

Cutlass通俗理解核心目的 核心目标: 消除矩阵乘法中对A, B矩阵的重复读取 核心: 内积变外积 -> 每个数据只会读第一次, 避免对A/B矩阵的重复读取 最内层外积, 可复用的地方也用外积, 体现就是MMA的warp level op 内积的问题: 如果每个thread负责C矩阵的一个元素C[m, n], 并且用下面这种方式循环(内积), 可以发现A, B矩阵会被重复...

Roofline Model and Flash Attention 计算强度: Flops/Byte => 反映每个字节的传输会产生多少计算 roofline model: 纵轴: 算力, 单位Flop/s 横轴: 计算强度I, 单位Flop/Byte 斜率: 内存带宽, 单位Byte/s 算力的上限决定了roofline model的屋顶 带宽的上限决定了roofline mode...

Tensor core MMA指令教程 参考 https://zhuanlan.zhihu.com/p/1892346599864238276 以mma.m8n8k4为例 A warp executing mma.m8n8k4 with .f16 floating point type will compute 4 MMA operations of shape .m8n8k4. 一个...

Deep GEMM解读terms _ss, _rs表示A/B矩阵的位置 math warp 用于计算的warp data warp 用于传输数据的warp st/ld store, load, 写回, 读出 barrier 同步信号 fence 防止乱序, 指令屏障 components wgmma 更大粒度的mma, warpgroup-level, 单次...