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

grpo cheat sheet https://huggingface.co/docs/trl/main/grpo_trainer GRPO(Group Relative Policy Optimization)的4个步骤: 其中1,2阶段相当于准备阶段, 3,4阶段相当真正的训练阶段 生成补完(Generating completions) AKA推理生成, 不带梯度 同一个输入...

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

内存序简记AKA: 控制使用Cache的行为 acquire: 间记: pull 保证一定刷新读缓存 comsume: 间记: 非pull 保证一定刷新写缓存

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

Ray UsageRay能方便使用多进程和分布式, 无需手动管理socket、线程池、消息队列等, 只需在函数或类前加上@ray.remote装饰器, 就能将其变成分布式任务或分布式对象。底层用了KV数据库自动管理这里metadata。 核心抽象 task: function call actor: 分布式对象, AKA object的集合 object: 分布式变量, AKA futur...

cutlass layout algebra: cheatsheet约定: 永远都是一维index的映射/变换, 只不过可以构造一些逻辑上的view/layout 逻辑空间: domain, 物理空间: co-domain 所以index的范围是相同的 column-major构建index layout[idx] = column-major地walk一个shape 对于一个shap...

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矩阵会被重复...