[webgpu] Restore MatMulNBits workgroup size for Phi-3.5 #23349
Conversation
This change restores the MatMulNBits workgroup size from (8, 8, 1) back to (16, 8, 1) to resolve a performance regression observed on Intel iGPUs during token generation (M=1). Signed-off-by: Jianhui Dai <[email protected]>
|
Please take a look before broader review. |
|
Do you have more specific information on how much this regresses on what GPU platforms? |
|
Phi-3.5 decoding performance on LNL decreased from approximately 28 tps to 26 tps. The detailed profiling results below. **workgroup_size-8x8x1
**workgroup_size-16x8x1
|
| @@ -583,7 +583,8 @@ Status MatMulNBits::ComputeInternal(onnxruntime::webgpu::ComputeContext& context | |||
| program.CacheHint("T_M" + std::to_string(tile_m) + "Subgroup" + std::to_string(use_subgroup)); | |||
| } else if (block_size == 32) { | |||
| components = 1; | |||
| constexpr uint32_t workgroup_size = 64; | |||
| // TODO: Tune the workgroup size when `M=1`. | |||
| constexpr uint32_t workgroup_size = 128; | |||
There was a problem hiding this comment.
Thanks for catching it. I changed it by accident. Maybe just restore it as before:
constexpr uint32_t workgroup_size = 128;
const uint32_t workgroup_y = N % 8 == 0 ? 8 : N % 4 == 0 ? 4
: 1;
There was a problem hiding this comment.
We tested various combinations of workgroup_size_x, workgroup_size_y, and output_number (e.g. (32, 8, 1), (32, 16, 1), (16, 16, 1) ...), with initial results showing a 5-15% performance improvement (M=1) compared to (16, 8, 1) , with results varying across devices.
This TODO tracks upcoming tuning.
Description
This change restores the MatMulNBits workgroup size from (8, 8, 1) back to (16, 8, 1) to resolve a performance regression observed on Intel iGPUs during token generation (M=1).
Motivation and Context
As above.