How to tune cutlass matmul kernels to approach cublasLt one?

[MARD1NO]

Recently I want to compare cublasLt kernel and cutlass kernel.

Actually cublasLt calls the cutlass matmul kernel, but I cannot get the same result, maybe I need to change some configuration in cutlass kernel.

I test the matmul+bias+gelu pattern, and I found there is an example: 12_gemm_bias_relu, I just simply change to use linearCombinationGELU epilogue.

The case is:

const int length_m = 4096;
const int length_n = 5120;
const int length_k = 4096;

And cublasLt's kernel name is: void cutlass::Kernel<cutlass_80_tensorop_f16_s16816gemm_gelu_f16_128x128_32x5_nn_align8>

I found the kernel naming rules in generators.py and change the configure as below:

using ShapeMMAThreadBlock = cutlass::gemm::GemmShape<128, 128, 32>; 
using ShapeMMAWarp = cutlass::gemm::GemmShape<64, 64, 32>; 
using ShapeMMAOp = cutlass::gemm::GemmShape<16, 8, 16>; 

using SwizzleThreadBlock = cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<8>;
constexpr int NumStages = 5; 

And I also found cublasLt use tanh to approximate GELU function:

So I change to use GELU_taylor epilogue:

using LinearCombinationFastGELU = LinearCombinationGeneric<GELU_taylor, ElementOutput_, Count, ElementAccumulator_,
                                                       ElementCompute_, Scale, Round, true>;

And it still have some gap:
cublasLt 855us vs cutlass 900us, and I also found the grid configuration is different. cublasLt is (320, 4, 2), cutlass is (320, 4, 1). And I wonder is there exists some configurations are not aligned with cublasLt？

[hwu36]

You are very good!

cublas has 2 in its grid.z which means it uses splitK. To do apple to apple comparison, you need to do the same.

Note, different version of cutlass and compiler will cause some performance difference.

[MARD1NO]

Finally, I found in cuda11.6, cutlass's performance is same as cublasLt one. In previous experiment, I use cuda11.7......

And I also compare with splitK version. origin 760us vs splitK2(the same as cublasLt) 800us

[hwu36]

Yeah. That is what you would see if you use 11.7. Our gtc22 fall talk mentioned the choice of the compiler.

[alohali]

@hwu36 hi, do you have notes about this? " Our gtc22 fall talk mentioned the choice of the compiler." we have no idea about compiler version choice but only use tot

[hwu36]

tot is good. next will be better. for gemm, 11.3+ is needed, just avoid 11.7

[alohali]

cool! thanks!