[CK_TILE] Update gfx11 FMHA forward kernel configs

[hyoon1]

Motivation

Tune gfx11 FMHA codegen to recover performance for mainly PSSK (padded seqlen_q/k) cases.
This tuning is based on heuristic search and improves performance in most tested shapes.
Performance should be evaluated on top of ROCm/rocm-libraries#5018 (required baseline).

Technical Details

• Updated gfx11 codegen heuristic choices for tile size and occupancy.
• Updated gfx11 pipeline selection:
  - Disabled the npad (f,f,f,f) qr entry because it was consistently slower than the pssk (t,t,f,f) path, and kept pssk enabled so npad cases are dispatched to the faster kernel path.`
• Kept gfx12 unchanged: with PSSK support from ROCm/rocm-libraries#4957, existing gfx12 config is already sufficient.
• Tuning rationale:
  • In some cases, higher kBlockPerCu lowers register pressure.
  • On RDNA, this generally aligns with better performance when waves_per_eu >= 6.

Test Plan

• test_ck_tile_fmha
• tile_example_fmha_fwd: tested this on gfx1100 and gfx1151
  ./build/bin/tile_example_fmha_fwd -prec=bf16 -mode={0/1} -b=1 -h=24 -d=128 -s={seqlen} -s_k={seqlen} -lse=0 -iperm={0/1} -operm={0/1}

Test Result

• TFLOPs by sequence length target: gfx1100 layout: bhsd
• mode: batch

SeqLen	Gain
1024	0.97x
4096	1.17x
8192	1.11x
12288	1.17x
16384	1.34x
20480	1.40x
24576	1.41x
27280	1.63x

• mode: group

SeqLen	Gain
1024	0.99x
4096	1.19x
8192	1.17x
12288	1.20x
16384	1.42x
20480	1.56x
24576	1.64x
27280	1.67x

• TFLOPs by sequence length target: gfx1151 layout: bshd
• mode: batch

Batch	Gain
1024	1.09x
4096	1.05x
8192	1.01x
12288	0.99x
16384	1.05x
20480	0.99x
24576	1.04x
27280	2.44x

• mode: group

Batch	Gain
1024	0.97x
4096	1.05x
8192	1.07x
12288	1.00x
16384	1.24x
20480	1.71x
24576	2.33x
27280	2.61x

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[ex-rzr]

What about other hdims? Is "f", "f", "f", "f" slower than "t", "t", "f", "f" for them too? (I know that padding for seqlen dimensions does not affect performance as much as padding for hdim, I'm just curious because it's really interesting observation)

Otherwise LGTM (btw, my approval does not make it possible to merge the PR because I'm not a code owner)

[hyoon1]

What about other hdims? Is "f", "f", "f", "f" slower than "t", "t", "f", "f" for them too? (I know that padding for seqlen dimensions does not affect performance as much as padding for hdim, I'm just curious because it's really interesting observation)

Otherwise LGTM (btw, my approval does not make it possible to merge the PR because I'm not a code owner)

It varied a bit depending on the GPU target and head dimension. In some cases ffff was slightly faster than ttff, but most of the time the performance was very similar and the difference wasn’t really noticeable.

Since the workloads I looked at mostly use hdim 64 or 128 and ttff already performs well there, it doesn’t seem worth adding extra kernel variants for now. Enabling only ttff should be sufficient.