XLA:数值检查通过;编译后的 HLO 中 cross_program_prefetch 无,S(3) 0 处,ROOT 是 fusion Pallas(out_type=pltpu.HBM):数值检查通过;编译后的 HLO 中 cross_program_prefetch 无,S(3) 0 处,ROOT 是 copy Pallas(out_type=ShapeDtypeStruct):数值检查通过;编译后的 HLO 中 cross_program_prefetch 无,S(3) 0 处,ROOT 是 custom-call ## 方法 1:每次调用并等到结果,固定同一个输入 XLA:231.5 µs Pallas(out_type=pltpu.HBM):275.3 µs Pallas(out_type=ShapeDtypeStruct):235.5 µs ## 方法 2:每次调用并等到结果,每次一个新的输入数组 XLA:232.1 µs Pallas(out_type=pltpu.HBM):271.1 µs Pallas(out_type=ShapeDtypeStruct):236.8 µs ## 方法 3:设备上的 LCC(KernelClock),每次一个新的输入数组 XLA,module:TensorCore 0 72476 个周期,TensorCore 1 36373 个周期 XLA,multiply_add_fusion:TensorCore 0 36256 个周期,TensorCore 1 36269 个周期 Pallas(out_type=pltpu.HBM),module:TensorCore 0 116322 个周期,TensorCore 1 80211 个周期 Pallas(out_type=pltpu.HBM),transform.1:TensorCore 0 43774 个周期,TensorCore 1 43798 个周期 Pallas(out_type=pltpu.HBM),copy:TensorCore 0 36256 个周期,TensorCore 1 36218 个周期 Pallas(out_type=ShapeDtypeStruct),module:TensorCore 0 80584 个周期,TensorCore 1 43964 个周期 Pallas(out_type=ShapeDtypeStruct),transform.1:TensorCore 0 43846 个周期,TensorCore 1 43866 个周期 ## 方法 4:在一个 jit 中用 fori_loop 重复 N 次,(T(32) − T(16)) / 16 XLA:每次 9.5 µs;循环版本的 HLO 中 cross_program_prefetch 无,S(3) 13 处,ROOT 是 copy-done Pallas(out_type=pltpu.HBM):每次 43.3 µs;循环版本的 HLO 中 cross_program_prefetch 无,S(3) 1 处,ROOT 是 get-tuple-element Pallas(out_type=ShapeDtypeStruct):每次 31.2 µs;循环版本的 HLO 中 cross_program_prefetch 无,S(3) 12 处,ROOT 是 copy-done ## 方法 5:方法 2 按 AB/BA 交替做 4 轮 第 1 轮(XLA、Pallas(out_type=pltpu.HBM)、Pallas(out_type=ShapeDtypeStruct)):XLA 224.5 µs,Pallas(out_type=pltpu.HBM) 266.9 µs,Pallas(out_type=ShapeDtypeStruct) 229.9 µs 第 2 轮(Pallas(out_type=ShapeDtypeStruct)、Pallas(out_type=pltpu.HBM)、XLA):XLA 223.5 µs,Pallas(out_type=pltpu.HBM) 266.1 µs,Pallas(out_type=ShapeDtypeStruct) 227.7 µs 第 3 轮(XLA、Pallas(out_type=pltpu.HBM)、Pallas(out_type=ShapeDtypeStruct)):XLA 220.9 µs,Pallas(out_type=pltpu.HBM) 266.3 µs,Pallas(out_type=ShapeDtypeStruct) 234.8 µs 第 4 轮(Pallas(out_type=ShapeDtypeStruct)、Pallas(out_type=pltpu.HBM)、XLA):XLA 223.9 µs,Pallas(out_type=pltpu.HBM) 268.9 µs,Pallas(out_type=ShapeDtypeStruct) 232.3 µs