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

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...

为什么需要swizzlee.g. 转置 ldmatrix 假设有4路bank, 列向能并行 转置e.g. 把如下矩阵转置。 120 1 2 3x x x x 此时4路bank只用了1路。bank conflict。 12340 x1 x2 x3 x ldmatrixldmatrix可以一次加载一个matrix。但是matrix的行和行之间不一定是连续的，所以会有bank confli...

cutlass swizzle简单理解 物理行号=逻辑行号, 物理列号=逻辑列号^逻辑组号 一个index/offset有明显的组别关系, e.g. 0b rr ccc b, 表示|行号bits|(组)列号bits|数据组bits| 用swizzle映射一个逻辑地址, 可以实现无bank conflict的smem访问: physical_addr = swizzle(logical_ad...

cutlass cute copy的本质本质: src_ptr和dst_ptr在step时用的逻辑index进行for循环, 通过layout映射到物理index再进行访存。 123456789101112template <class SrcEngine, class SrcLayout, class DstEngine, class DstLayout>CU...

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

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. 一个...