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

为什么需要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实现flash attentionflash attention自顶向下(虽然我学cutlass是自底向上学的但是感觉快速上手应该自顶向下学)。因为有了cutlass cute用户就可以方便的实现一些功能了, 即一些cuda编程的范式: cuda程序范式: global mem -> share mem -> reg -> compute blo...