首页
学习
活动
专区
工具
TVP
发布
精选内容/技术社群/优惠产品,尽在小程序
立即前往

《PytorchConference2023 翻译系列》7-深入探索CUTLASS:如何充分利用Tensor Cores​​

我们支持多种Epilogue模式以及在NVIDIA GPU上找到的所有数据计算类型。我们最近发布了一个Python接口,我稍后会详细介绍它。...Epilogue(收尾)是主循环之后的一系列操作,用于处理主循环的输出结果。它可以执行各种操作,如对结果进行修正、缩放、舍入等。Epilogue的目的是将主循环的输出转换为最终的矩阵乘累加结果。...CUTLASS提供了不同类型的Epilogue,可以根据具体需求选择适当的Epilogue类型。...::collective::DefaultEpilogue cutlass::epilogue::collective::Epilogue Tiled (MMA and Copy) cute::TiledMma...这将允许用户使用基本的epilogue单元来开发复杂的epilogue。它是一组小的计算、加载和存储操作,可以生成常见或自定义的epilogue

57310
您找到你想要的搜索结果了吗?
是的
没有找到

大幅优化推理过程,字节高性能Transformer推理库获IPDPS 2023最佳论文奖

因为 softmax 需要对整行数据做归约,但因为共享内存大小的限制,一个 threadblock 内不能容纳整行数据,同时 threadblock 间的通信很低效,所以不能仅在 Q x K 的 epilogue...把 softmax 拆分成三步计算,分别 fuse 到 Q x K 的 epilogue 中, QK x V 的 prologue 中,以及中间再添加一个轻量的 kernel 做规约。...分为三步计算,大部分计算 fuse 到前后的 GEMM kernel 中 算法步骤: partial reduction:Q x K 的 epilogue 中,每个 threadblock 内部规约,计算出...GEMM & add-bias & GELU fusion 通过 CUTLASS fuse epilogue 的方式,把矩阵乘后的 add-bias 操作和 GELU activation 操作 fuse...拆解各优化手段对性能的影响如下: add-bias & LayerNorm fusion 可以提高性能 3.2% 将 add-bias & GELU fuse 到 GEMM epilogue 可以进一步提高

54310
领券