Compilation failure with modified tuning parameters on CK Tile GEMM

[zjing14]

composable_kernel/example/ck_tile/03_gemm/gemm_mem_pipeline.cpp

Lines 21 to 31 in 365f39a

	constexpr ck_tile::index_t M_Tile = 128;
	constexpr ck_tile::index_t N_Tile = 128;
	constexpr ck_tile::index_t K_Tile = 32;
	constexpr ck_tile::index_t M_Warp = 2;
	constexpr ck_tile::index_t N_Warp = 2;
	constexpr ck_tile::index_t K_Warp = 1;
	constexpr ck_tile::index_t M_Warp_Tile = 32;
	constexpr ck_tile::index_t N_Warp_Tile = 32;
	constexpr ck_tile::index_t K_Warp_Tile = 8;

Got the compilation failure with the following changes

    constexpr ck_tile::index_t M_Tile = 16;
    constexpr ck_tile::index_t N_Tile = 16;
    constexpr ck_tile::index_t K_Tile = 64;

    constexpr ck_tile::index_t M_Warp = 1;
    constexpr ck_tile::index_t N_Warp = 1;
    constexpr ck_tile::index_t K_Warp = 1;

    constexpr ck_tile::index_t M_Warp_Tile = 16;
    constexpr ck_tile::index_t N_Warp_Tile = 16;
    constexpr ck_tile::index_t K_Warp_Tile = 8;

[zjing14]

@aosewski @carlushuang

[illsilin]

@zjing14 Which compiler version did you use and what was the error message?

[poyenc]

I think we only have 16x16x16 warp gemm for now (v_mfma_f32_16x16x16f16/v_mfma_f32_16x16x16bf16_1k). And 16x16x8 warp gemm is only for bf16 (v_mfma_f32_16x16x8bf16), but we haven't added it in ck_tile.

[ThomasNing]

Poyen is right. Just add a little bit, alternatively, we could have the v_mfma_f32_32x32x8_f16: 32x32x8 warp tile and v_mfma_f32_16x16x4_4b_f16: 16x16x4 for 4 matrix batches. We could add the first instruction set support to the ck_tile gemm first.

[aosewski]

@zjing14 If you would change k warp tile to 16: K_Warp_Tile = 16; and change in example to use universal gemm pipeline as in gemm_basic.cpp example then it should work. This example just doesn't have yet enabled universal gemm pipeline as of the date of writing. Will soon change all examples to use universal gemm pipeline when we will be sure that it works best.

[zjing14]

@poyenc @ThomasNing @illsilin @aosewski Thank you all very much. I will give a try with new changes.

[darren-amd]

Hi @zjing14,

Have you gotten a chance to give it a try?

[zjing14]

@darren-amd Tried 16x16 Tile with ../example/ck_tile/03_gemm/gemm_basic.cpp, still got a compilation failure. What did I miss?

    constexpr ck_tile::index_t M_Tile = 16;
    constexpr ck_tile::index_t N_Tile = 16;
    constexpr ck_tile::index_t K_Tile = 32;

    constexpr ck_tile::index_t M_Warp = 1;
    constexpr ck_tile::index_t N_Warp = 1;
    constexpr ck_tile::index_t K_Warp = 1;

    constexpr ck_tile::index_t M_Warp_Tile = 16;
    constexpr ck_tile::index_t N_Warp_Tile = 16;
    constexpr ck_tile::index_t K_Warp_Tile = 8;

[darren-amd]

Hey @zjing14,

Could you try changing the last line to 16, as in K_Warp_Tile = 16; and give that a try? Thanks!

[zjing14]

@darren-amd Oh, I see. It works, thanks.

[zjing14]

BTW, Got a failure with example/ck_tile/03_gemm/gemm_mem_pipeline.cpp