What should I do when I want fuse bias add in Hopper Gemm?

[MARD1NO]

I just use the 54_hopper_fp8_warp_specialized_gemm.cu example and it use cutlass::epilogue::fusion::ScaledLinCombPerRowBiasEltActAmaxAux, (For my case I do not need the amax and aux output, so I change to use ScaledLinCombPerRowBiasEltAct this one). Without bias add, its result is correct, but when I pass bias_ptr to it, it got wrong result.

And I also check some cutlass test example, I found they use Sm90ScaledLinCombPerRowBiasEltAct, is this needed? And I'm not very clear how to pass arguments like scale, bias_ptr, for this new FusionCallback, can you offer some example?

Thanks~:D

emmm I found the example in sm90_callbacks_tma_warpspecialized.hpp

[MARD1NO]

Just make a mistake, row bias pattern is add (M, 1) shape bias, and I want to add (1, N) shape bias, wait for Sm90ScaledLinCombPerColBiasEltAct :D

[hwu36]

@apuaaChen to comment on column bias fusion