[llvm] [AMDGPU] Fix opsel for scaled MFMA operations (PR #140183)

[Shilei Tian via llvm-commits]

================
@@ -445,23 +445,23 @@ v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], 33, 9
 // ERR: :[[@LINE+1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 cbsz:3 blgp:1
 
-// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x00,0xac,0xd3,0x18,0x33,0x02,0x08,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]
+// GFX950: v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[20:23], v24, v25 op_sel:[0,1,0] op_sel_hi:[0,0,0] cbsz:3 blgp:1 ; encoding: [0x00,0x10,0xac,0xd3,0x18,0x33,0x02,0x00,0x00,0x0b,0xad,0xd3,0x04,0x19,0x52,0x24]
----------------
shiltian wrote:

A side question: SP3 says this should be:

```
v_mfma_scale_f32_16x16x128_f8f6f4  v[0:3], v[4:11], v[12:19], v[20:23], v24, v25 op_sel:[0,1] op_sel_hi:[0,0] cbsz:3 abid:1 blgp:1
```

Why do we not emit `abid:1` here?

https://github.com/llvm/llvm-project/pull/140183