[Kernel][Triton][AMD] Change default block size for triton_scaled_mm to 128 for 3-5x speedup

[rasmith]

Changed default block-size for triton_scaled_mm to 128x128x128 from 32x32x32 for better performance. This results in roughly 3-5x speedup.

python benchmarks/benchmark_latency.py --dtype bfloat16 --enable-chunked-prefill False --load-format dummy --batch-size 64 --num-iters-warmup 2 --num-iters 5 --input-len 2048 --output-len 128 --model /models/Phi-3-medium-128k-instruct-quantized.w8a8/

Before:

Avg latency: 14.48 seconds

After:

Avg latency: 5.52 seconds

python benchmarks/benchmark_throughput.py --dtype bfloat16 --enable-chunked-prefill False --load-format dummy --input-len 2048 --output-len 128 --model /models/Phi-3-medium-128k-instruct-quantized.w8a8/

Before:

Throughput: 10269.32 tok/s

After:

Throughput: 31150.8 tokens/s

[github-actions]

👋 Hi! Thank you for contributing to the vLLM project.
Just a reminder: PRs would not trigger full CI run by default. Instead, it would only run fastcheck CI which starts running only a small and essential subset of CI tests to quickly catch errors. You can run other CI tests on top of those by going to your fastcheck build on Buildkite UI (linked in the PR checks section) and unblock them. If you do not have permission to unblock, ping simon-mo or khluu to add you in our Buildkite org.

Once the PR is approved and ready to go, your PR reviewer(s) can run CI to test the changes comprehensively before merging.

To run CI, PR reviewers can do one of these:

• Add ready label to the PR
• Enable auto-merge.

🚀

[mgoin]

This is an impressive improvement! Could you also show comparisons for equal input len/output len workloads, preferably with low batchsize? This could regress the TPOT for small decode batches.

It seems there is no tuning for this kernel at the moment, so maybe this could benefit from a simple heuristic for the extreme problem sizes or a few @triton.autotune configs for the blocksizes.