AMDGPU : Fix mqsad_u32_u8 instruction incorrect data type.

[wdng]

Summary

1. Ouput and the third parameter should be v4i32.
2. Add a 128 register class to hand the v4i32 case.

[arsenm]

A lot more tests for the operand combinations are needed. There should be a test where every operand is a VGPR, SGPR, inline integer immediate, inline FP immediate, and non-inline constant

include/llvm/IR/Intrinsics.td

152 ↗

(On Diff #68638)

You don't need this

lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

272–276 ↗

(On Diff #68638)

You don't need SSrc128

lib/Target/AMDGPU/SIRegisterInfo.h

118–119 ↗

(On Diff #68638)

Indentation. This probably isn't needed to be handled since it's starting out as isAllocatable = 0

[wdng]

In D23700#520337, @arsenm wrote:

A lot more tests for the operand combinations are needed.

Do you mean to use some combination like v2i64?

[arsenm]

In D23700#521053, @wdng wrote:

In D23700#520337, @arsenm wrote:

A lot more tests for the operand combinations are needed.

Do you mean to use some combination like v2i64?

No, I mean the kinds of values that I listed

[wdng]

In D23700#521071, @arsenm wrote:

In D23700#521053, @wdng wrote:

In D23700#520337, @arsenm wrote:

A lot more tests for the operand combinations are needed.

Do you mean to use some combination like v2i64?

No, I mean the kinds of values that I listed

Where are those values?

[wdng]

Update LIT tests to make sure using operand as a VGPR, SGPR, inline integer immediate, inline FP immediate, and non-inline constant.

[arsenm]

I think we need some execution tests to make sure the inline immediates and quad SGPR really work

test/CodeGen/AMDGPU/llvm.amdgcn.mqsad.u32.u8.ll

25–26 ↗

(On Diff #68887)

This is broken. Storing the constant to undef is not useful. The high bits also need to be 0, and the constant vector should be the direct operand to the call

34–38 ↗

(On Diff #68887)

None of these are inline fp immediates, and you don't want to do an fp cast to integer. You want the literal bit value for the FP immediate

46–56 ↗

(On Diff #68887)

Nothing here is constraining the input to be a VGPR. The simplest way to get an SGPR is to use a kernel argument value. To get a VGPR you should use a load from addrspace(1)* or use inline asm

57 ↗

(On Diff #68887)

Space before = (same for the other tests)

[wdng]

Update LIT tests.

[wdng]

New patch has fixed those issues.

[arsenm]

lib/Target/AMDGPU/SIRegisterInfo.td

348–350 ↗

(On Diff #68928)

We don't need this since the operand is required to be VReg_128

[wdng]

lib/Target/AMDGPU/SIRegisterInfo.td

348–350 ↗

(On Diff #68928)

Are you sure that we don' t need this? Since VS_128 is used a lot by other definitions...

[wdng]

Remove 128-bit operand type definition and use VReg_128 directly.

[arsenm]

lib/Target/AMDGPU/SIInstrInfo.td

1123 ↗

(On Diff #69395)

This should not be VCSrc_128, because an immediate is not allowed

[wdng]

Use VSrc to replace VCSrc to take care of the immediate operand.

[arsenm]

Because this is a new width operand, there should be some assembler tests for it

lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

276–278 ↗

(On Diff #69410)

You should not need this

lib/Target/AMDGPU/SIRegisterInfo.td

403 ↗

(On Diff #69410)

This should just be RegisterOperand

404 ↗

(On Diff #69410)

I don't think you need the custom matcher

[wdng]

1. Changes based on Matt's comments.
2. Add assembler tests.

[arsenm]

The tests should be merged into the complete vop3 tests

[wdng]

Merge assembler tests.

[arsenm]

test/MC/AMDGPU/vop3.s

1–2 ↗

(On Diff #70225)

Only the NOSI one should be checking the stderr. The -mcpu=SI shouldalso be removed. This breaks all of the SICI check lines

[wdng]

Modified check-lines based on Matt's comments.

[arsenm]

test/MC/AMDGPU/vop3.s

363 ↗

(On Diff #70439)

CIVI doesn't make sense, and you don't have such a prefix. You need a separate CI and VI check line for this

[wdng]

Separate CI and VI check lines.

[arsenm]

test/MC/AMDGPU/vop3.s

363 ↗

(On Diff #70537)

The check prefix is SICI but you are only using CI here

[wdng]

test/MC/AMDGPU/vop3.s

363 ↗

(On Diff #70537)

But v_mqsad won't work on SI, so I separate CI and VI for this test. For SI, NOSI is used to check the stderr.

[arsenm]

test/MC/AMDGPU/vop3.s

363 ↗

(On Diff #70537)

That doesn't matter, you just need to use the same check prefix. CI isn't a check prefix so it isn't checking anything

[wdng]

test/MC/AMDGPU/vop3.s

363 ↗

(On Diff #70537)

If I change it to:
// SICI: v_mqsad_u32_u8 v[0:3], s[2:3], v4, v[0:3] ; encoding: [0x00,0x00,0xe7,0xd1,0x02,0x08,0x02,0x04]

It failed with the following error:

/llvm/test/MC/AMDGPU/vop3.s:362:1: error: instruction not supported on this GPU
v_mqsad_u32_u8 v[0:3], s[2:3], v4, v[0:3]
^
/llvm/test/MC/AMDGPU/vop3.s:363:10: error: expected string not found in input
// SICI: v_mqsad_u32_u8 v[0:3], s[2:3], v4, v[0:3] ; encoding: [0x00,0x00,0xe7,0xd1,0x02,0x08,0x02,0x04]

[arsenm]

test/MC/AMDGPU/vop3.s

363 ↗

(On Diff #70537)

You need to use a CI specific checkline

[wdng]

Add a separate CI check-line.

[arsenm]

test/MC/AMDGPU/vop3.s

4 ↗

(On Diff #70872)

tonga is not CI

[wdng]

Set mcpu=hawaii for CI.

[arsenm]

LGTM

llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td

def int_amdgcn_qsad_pk_u16_u8 :
   Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty], [IntrNoMem]>;
 
 def int_amdgcn_mqsad_pk_u16_u8 :
   GCCBuiltin<"__builtin_amdgcn_mqsad_pk_u16_u8">,
   Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty], [IntrNoMem]>;
 
 def int_amdgcn_mqsad_u32_u8 :
   GCCBuiltin<"__builtin_amdgcn_mqsad_u32_u8">,
-  Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
+  Intrinsic<[llvm_v4i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_v4i32_ty], [IntrNoMem]>;
 
 def int_amdgcn_cvt_pk_u8_f32 :
   GCCBuiltin<"__builtin_amdgcn_cvt_pk_u8_f32">,
   Intrinsic<[llvm_i32_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 
 def int_amdgcn_icmp :
   Intrinsic<[llvm_i64_ty], [llvm_anyint_ty, LLVMMatchType<0>, llvm_i32_ty],
             [IntrNoMem, IntrConvergent]>;

llvm/trunk/lib/Target/AMDGPU/CIInstructions.td

 defm V_MQSAD_U16_U8 : VOP3Inst <vop3<0x172>, "v_mqsad_u16_u8",
   VOP_I32_I32_I32
 >;
 
 defm V_QSAD_PK_U16_U8 : VOP3Inst <vop3<0x172, 0x1e5>, "v_qsad_pk_u16_u8",
   VOP_I64_I64_I32_I64, int_amdgcn_qsad_pk_u16_u8>;
 
 defm V_MQSAD_U32_U8 : VOP3Inst <vop3<0x174, 0x1e7>, "v_mqsad_u32_u8",
-  VOP_I32_I32_I32_I32, int_amdgcn_mqsad_u32_u8>;
+  VOP_V4I32_I64_I32_V4I32, int_amdgcn_mqsad_u32_u8>;
 
 let isCommutable = 1 in {
 defm V_MAD_U64_U32 : VOP3Inst <vop3<0x176>, "v_mad_u64_u32",
   VOP_I64_I32_I32_I64
 >;
 
 // XXX - Does this set VCC?
 defm V_MAD_I64_I32 : VOP3Inst <vop3<0x177>, "v_mad_i64_i32",

llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td

!if (!eq(Src0.Value, untyped.Value),      0,
          !if (!eq(Src2.Value, untyped.Value), 2,   // VOP2
                                               3))); // VOP3
 }
 
 // Returns the register class to use for the destination of VOP[123C]
 // instructions for the given VT.
 class getVALUDstForVT<ValueType VT> {
   RegisterOperand ret = !if(!eq(VT.Size, 32), VOPDstOperand<VGPR_32>,
+                          !if(!eq(VT.Size, 128), VOPDstOperand<VReg_128>,
                             !if(!eq(VT.Size, 64), VOPDstOperand<VReg_64>,
                               !if(!eq(VT.Size, 16), VOPDstOperand<VGPR_32>,
-                            VOPDstOperand<SReg_64>))); // else VT == i1
+                              VOPDstOperand<SReg_64>)))); // else VT == i1
 }
 
 // Returns the register class to use for source 0 of VOP[12C]
 // instructions for the given VT.
 class getVOPSrc0ForVT<ValueType VT> {
   bit isFP = !if(!eq(VT.Value, f16.Value), 1,
              !if(!eq(VT.Value, f32.Value), 1,
              !if(!eq(VT.Value, f64.Value), 1,
              0)));
-  RegisterOperand ret = !if(isFP, 
+  RegisterOperand ret = !if(isFP,
                             !if(!eq(VT.Size, 64), VSrc_f64, VSrc_f32),
                             !if(!eq(VT.Size, 64), VSrc_b64, VSrc_b32));
 }
 
 // Returns the vreg register class to use for source operand given VT
 class getVregSrcForVT<ValueType VT> {
-  RegisterClass ret = !if(!eq(VT.Size, 64), VReg_64, VGPR_32);
+  RegisterClass ret = !if(!eq(VT.Size, 128), VReg_128,
+                        !if(!eq(VT.Size, 64), VReg_64, VGPR_32));
 }
 
 
 // Returns the register class to use for sources of VOP3 instructions for the
 // given VT.
 class getVOP3SrcForVT<ValueType VT> {
   bit isFP = !if(!eq(VT.Value, f16.Value), 1,
              !if(!eq(VT.Value, f32.Value), 1,
              !if(!eq(VT.Value, f64.Value), 1,
              0)));
   RegisterOperand ret =
+  !if(!eq(VT.Size, 128),
+      VSrc_128,
     !if(!eq(VT.Size, 64),
         !if(isFP,
             VCSrc_f64,
             VCSrc_b64),
         !if(!eq(VT.Value, i1.Value),
             SCSrc_b64,
             !if(isFP,
                 VCSrc_f32,
                 VCSrc_b32)
          )
-      );
+     )
 }
 
 // Returns 1 if the source arguments have modifiers, 0 if they do not.
 // XXX - do f16 instructions?
 class hasModifiers<ValueType SrcVT> {
   bit ret =
     !if(!eq(SrcVT.Value, f32.Value), 1,
     !if(!eq(SrcVT.Value, f64.Value), 1,

 class getInsSDWA <RegisterClass Src0RC, RegisterClass Src1RC, int NumSrcArgs,
                   bit HasFloatModifiers, Operand Src0Mod, Operand Src1Mod,
                   ValueType DstVT> {
 
   dag ret = !if(!eq(NumSrcArgs, 0),
                // VOP1 without input operands (V_NOP)
                (ins),
             !if(!eq(NumSrcArgs, 1),
-                !if(HasFloatModifiers, 
+                !if(HasFloatModifiers,
                     // VOP1_SDWA with float modifiers
                     (ins Src0Mod:$src0_fmodifiers, Src0RC:$src0,
                          clampmod:$clamp, dst_sel:$dst_sel, dst_unused:$dst_unused,
                          src0_sel:$src0_sel),
                     // VOP1_SDWA with int modifiers
                     (ins Src0Mod:$src0_imodifiers, Src0RC:$src0,
                          clampmod:$clamp, dst_sel:$dst_sel, dst_unused:$dst_unused,
                          src0_sel:$src0_sel))
               /* NumSrcArgs == 2 */,
               !if(HasFloatModifiers,
                   !if(!eq(DstVT.Size, 1),
                       // VOPC_SDWA with float modifiers
                       (ins Src0Mod:$src0_fmodifiers, Src0RC:$src0,
                            Src1Mod:$src1_fmodifiers, Src1RC:$src1,
                            clampmod:$clamp, src0_sel:$src0_sel, src1_sel:$src1_sel),
                       // VOP2_SDWA or VOPC_SDWA with float modifiers
                       (ins Src0Mod:$src0_fmodifiers, Src0RC:$src0,
                            Src1Mod:$src1_fmodifiers, Src1RC:$src1,
                            clampmod:$clamp, dst_sel:$dst_sel, dst_unused:$dst_unused,
                            src0_sel:$src0_sel, src1_sel:$src1_sel)),
-              
                   !if(!eq(DstVT.Size, 1),
                       // VOPC_SDWA with int modifiers
                       (ins Src0Mod:$src0_imodifiers, Src0RC:$src0,
                            Src1Mod:$src1_imodifiers, Src1RC:$src1,
                            clampmod:$clamp, src0_sel:$src0_sel, src1_sel:$src1_sel),
                       // VOP2_SDWA or VOPC_SDWA with int modifiers
                       (ins Src0Mod:$src0_imodifiers, Src0RC:$src0,
                            Src1Mod:$src1_imodifiers, Src1RC:$src1,
def VOP_MAC : VOPProfile <[f32, f32, f32, f32]> {
   let AsmDPP = getAsmDPP<1, 2, HasModifiers, f32>.ret;
   let AsmSDWA = getAsmSDWA<1, 2, HasModifiers, f32>.ret;
 }
 def VOP_F64_F64_F64_F64 : VOPProfile <[f64, f64, f64, f64]>;
 def VOP_I32_I32_I32_I32 : VOPProfile <[i32, i32, i32, i32]>;
 def VOP_I64_I32_I32_I64 : VOPProfile <[i64, i32, i32, i64]>;
 def VOP_I32_F32_I32_I32 : VOPProfile <[i32, f32, i32, i32]>;
 def VOP_I64_I64_I32_I64 : VOPProfile <[i64, i64, i32, i64]>;
+def VOP_V4I32_I64_I32_V4I32 : VOPProfile <[v4i32, i64, i32, v4i32]>;
 
 // This class is used only with VOPC instructions. Use $sdst for out operand
 class SIInstAlias <string asm, Instruction inst, VOPProfile p,
                    string VariantName = ""> :
     InstAlias <asm, (inst)>, PredicateControl {
 
   field bit isCompare;
   field bit isCommutable;

llvm/trunk/lib/Target/AMDGPU/SIRegisterInfo.td

 defm SCSrc : RegInlineOperand<"SReg", "SCSrc"> ;
 
 //===----------------------------------------------------------------------===//
 //  VSrc_* Operands with an SGPR, VGPR or a 32-bit immediate
 //===----------------------------------------------------------------------===//
 
 defm VSrc : RegImmOperand<"VS", "VSrc">;
 
+def VSrc_128 : RegisterOperand<VReg_128>;
+
 //===----------------------------------------------------------------------===//
 //  VSrc_* Operands with an VGPR
 //===----------------------------------------------------------------------===//
 
 // This is for operands with the enum(9), VSrc encoding restriction,
 // but only allows VGPRs.
 def VRegSrc_32 : RegisterOperand<VGPR_32> {
   //let ParserMatchClass = RegImmMatcher<"VRegSrc32">;
   let DecoderMethod = "DecodeVS_32RegisterClass";
 }
 
 //===----------------------------------------------------------------------===//
 //  VCSrc_* Operands with an SGPR, VGPR or an inline constant
 //===----------------------------------------------------------------------===//
 
 defm VCSrc : RegInlineOperand<"VS", "VCSrc">;

llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.mqsad.u32.u8.ll

 ; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
 ; RUN: llc -march=amdgcn -mcpu=fiji -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
 
-declare i32 @llvm.amdgcn.mqsad.u32.u8(i32, i32, i32) #0
+declare <4 x i32> @llvm.amdgcn.mqsad.u32.u8(i64, i32, <4 x i32>) #0
 
-; GCN-LABEL: {{^}}v_mqsad_u32_u8:
-; GCN: v_mqsad_u32_u8 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
-define void @v_mqsad_u32_u8(i32 addrspace(1)* %out, i32 %src) {
-  %result= call i32 @llvm.amdgcn.mqsad.u32.u8(i32 %src, i32 100, i32 100) #0
-  store i32 %result, i32 addrspace(1)* %out, align 4
+; GCN-LABEL: {{^}}v_mqsad_u32_u8_use_non_inline_constant:
+; GCN: v_mqsad_u32_u8 v[{{[0-9]+:[0-9]+}}], s[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v[{{[0-9]+:[0-9]+}}]
+define void @v_mqsad_u32_u8_use_non_inline_constant(<4 x i32> addrspace(1)* %out, i64 %src) {
+  %result = call <4 x i32> @llvm.amdgcn.mqsad.u32.u8(i64 %src, i32 100, <4 x i32> <i32 100, i32 100, i32 100, i32 100>) #0
+  store <4 x i32> %result, <4 x i32> addrspace(1)* %out, align 4
   ret void
 }
 
 ; GCN-LABEL: {{^}}v_mqsad_u32_u8_non_immediate:
-; GCN: v_mqsad_u32_u8 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
-define void @v_mqsad_u32_u8_non_immediate(i32 addrspace(1)* %out, i32 %src, i32 %a, i32 %b) {
-  %result= call i32 @llvm.amdgcn.mqsad.u32.u8(i32 %src, i32 %a, i32 %b) #0
-  store i32 %result, i32 addrspace(1)* %out, align 4
+; GCN: v_mqsad_u32_u8 v[{{[0-9]+:[0-9]+}}], s[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v[{{[0-9]+:[0-9]+}}]
+define void @v_mqsad_u32_u8_non_immediate(<4 x i32> addrspace(1)* %out, i64 %src, i32 %a, <4 x i32> %b) {
+  %result = call <4 x i32> @llvm.amdgcn.mqsad.u32.u8(i64 %src, i32 %a, <4 x i32> %b) #0
+  store <4 x i32> %result, <4 x i32> addrspace(1)* %out, align 4
+  ret void
+}
+
+; GCN-LABEL: {{^}}v_mqsad_u32_u8_inline_integer_immediate:
+; GCN: v_mqsad_u32_u8 v[{{[0-9]+:[0-9]+}}], s[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v[{{[0-9]+:[0-9]+}}]
+define void @v_mqsad_u32_u8_inline_integer_immediate(<4 x i32> addrspace(1)* %out, i64 %src, i32 %a) {
+  %result = call <4 x i32> @llvm.amdgcn.mqsad.u32.u8(i64 %src, i32 %a, <4 x i32> <i32 10, i32 20, i32 30, i32 40>) #0
+  store <4 x i32> %result, <4 x i32> addrspace(1)* %out, align 4
+  ret void
+}
+
+; GCN-LABEL: {{^}}v_mqsad_u32_u8_inline_fp_immediate:
+; GCN: v_mqsad_u32_u8 v[{{[0-9]+:[0-9]+}}], s[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v[{{[0-9]+:[0-9]+}}]
+define void @v_mqsad_u32_u8_inline_fp_immediate(<4 x i32> addrspace(1)* %out, i64 %src, i32 %a) {
+  %result = call <4 x i32> @llvm.amdgcn.mqsad.u32.u8(i64 %src, i32 %a, <4 x i32> <i32 1065353216, i32 0, i32 0, i32 0>) #0
+  store <4 x i32> %result, <4 x i32> addrspace(1)* %out, align 4
+  ret void
+}
+
+; GCN-LABEL: {{^}}v_mqsad_u32_u8_use_sgpr_vgpr:
+; GCN: v_mqsad_u32_u8 v[{{[0-9]+:[0-9]+}}], s[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v[{{[0-9]+:[0-9]+}}]
+define void @v_mqsad_u32_u8_use_sgpr_vgpr(<4 x i32> addrspace(1)* %out, i64 %src, i32 %a, <4 x i32> addrspace(1)* %input) {
+  %in = load <4 x i32>, <4 x i32> addrspace(1) * %input
+
+  %result = call <4 x i32> @llvm.amdgcn.mqsad.u32.u8(i64 %src, i32 %a, <4 x i32> %in) #0
+  store <4 x i32> %result, <4 x i32> addrspace(1)* %out, align 4
   ret void
 }
 
 attributes #0 = { nounwind readnone }

llvm/trunk/test/MC/AMDGPU/vop3-errs.s

 // RUN: not llvm-mc -arch=amdgcn -show-encoding %s 2>&1 | FileCheck %s
 // RUN: not llvm-mc -arch=amdgcn -mcpu=SI -show-encoding %s 2>&1 | FileCheck %s
 
 v_add_f32_e64 v0, v1
 // CHECK: error: too few operands for instruction
 
 v_div_scale_f32  v24, vcc, v22, 1.1, v22
 // CHECK: error: invalid operand for instruction
+
+v_mqsad_u32_u8 v[0:3], s[2:3], v4, v[0:3]
+// CHECK: error: instruction not supported on this GPU

llvm/trunk/test/MC/AMDGPU/vop3.s

-// RUN: llvm-mc -arch=amdgcn -show-encoding %s | FileCheck %s --check-prefix=SICI
-// RUN: llvm-mc -arch=amdgcn -mcpu=SI -show-encoding %s | FileCheck %s --check-prefix=SICI
+// RUN: not llvm-mc -arch=amdgcn -show-encoding %s 2>&1 | FileCheck %s --check-prefix=NOSI
+// RUN: not llvm-mc -arch=amdgcn -show-encoding %s | FileCheck %s --check-prefix=SICI
+
+// RUN: llvm-mc -arch=amdgcn -mcpu=hawaii -show-encoding %s | FileCheck %s --check-prefix=CI
 // RUN: not llvm-mc -arch=amdgcn -mcpu=tonga -show-encoding %s | FileCheck %s --check-prefix=VI
 // RUN: not llvm-mc -arch=amdgcn -mcpu=tonga -show-encoding %s 2>&1 | FileCheck %s -check-prefix=NOVI
 
 
 //===----------------------------------------------------------------------===//
 // VOPC Instructions
 //===----------------------------------------------------------------------===//
 

 
 v_div_scale_f32 v24, vcc, v22, v22, 0xc0000000
 // SICI: v_div_scale_f32 v24, vcc, v22, v22, -2.0 ; encoding: [0x18,0x6a,0xda,0xd2,0x16,0x2d,0xd6,0x03]
 // VI:   v_div_scale_f32 v24, vcc, v22, v22, -2.0 ; encoding: [0x18,0x6a,0xe0,0xd1,0x16,0x2d,0xd6,0x03]
 
 v_mad_f32 v9, 0.5, v5, -v8
 // SICI: v_mad_f32 v9, 0.5, v5, -v8      ; encoding: [0x09,0x00,0x82,0xd2,0xf0,0x0a,0x22,0x84]
 // VI:   v_mad_f32 v9, 0.5, v5, -v8      ; encoding: [0x09,0x00,0xc1,0xd1,0xf0,0x0a,0x22,0x84]
+
+v_mqsad_u32_u8 v[0:3], s[2:3], v4, v[0:3]
+// CI: v_mqsad_u32_u8 v[0:3], s[2:3], v4, v[0:3] ; encoding: [0x00,0x00,0xe8,0xd2,0x02,0x08,0x02,0x04]
+// VI: v_mqsad_u32_u8 v[0:3], s[2:3], v4, v[0:3] ; encoding: [0x00,0x00,0xe7,0xd1,0x02,0x08,0x02,0x04]
+// NOSI: error: instruction not supported on this GPU