[llvm-branch-commits] [llvm] AMDGPU: Use RegisterOperand for MIMG class data operands (PR #157215)

[Matt Arsenault via llvm-branch-commits]

https://github.com/arsenm created 
https://github.com/llvm/llvm-project/pull/157215

Avoid using getLdStRegisterOperand hidden at the bottom
of the class hierarchy.

>From a72a09db9aa9a037860dab5a50a7ec39629439ee Mon Sep 17 00:00:00 2001
From: Matt Arsenault <matthew.arsena...@amd.com>
Date: Wed, 3 Sep 2025 15:31:08 +0900
Subject: [PATCH] AMDGPU: Use RegisterOperand for MIMG class data operands

Avoid using getLdStRegisterOperand hidden at the bottom
of the class hierarchy.
---
 llvm/lib/Target/AMDGPU/MIMGInstructions.td    | 134 +++----
 llvm/lib/Target/AMDGPU/SIInstrInfo.td         |  35 +-
 llvm/lib/Target/AMDGPU/SIRegisterInfo.td      |   2 +-
 llvm/test/MC/AMDGPU/gfx90a_ldst_acc.s         | 362 +++++++++---------
 .../MC/AMDGPU/misaligned-vgpr-tuples-err.s    |  20 +-
 5 files changed, 280 insertions(+), 273 deletions(-)

diff --git a/llvm/lib/Target/AMDGPU/MIMGInstructions.td 
b/llvm/lib/Target/AMDGPU/MIMGInstructions.td
index 5d9e4401931f4..bf787b230067d 100644
--- a/llvm/lib/Target/AMDGPU/MIMGInstructions.td
+++ b/llvm/lib/Target/AMDGPU/MIMGInstructions.td
@@ -420,7 +420,7 @@ class VSAMPLE_gfx12<int op, dag outs, int num_addrs, string 
dns="",
 }
 
 class MIMG_NoSampler_Helper <mimgopc op, string asm,
-                             RegisterClass dst_rc,
+                             RegisterOperand dst_rc,
                              RegisterClass addr_rc,
                              string dns="">
   : MIMG_gfx6789 <op.GFX10M, (outs dst_rc:$vdata), dns> {
@@ -433,10 +433,10 @@ class MIMG_NoSampler_Helper <mimgopc op, string asm,
 }
 
 class MIMG_NoSampler_Helper_gfx90a <mimgopc op, string asm,
-                                    RegisterClass dst_rc,
+                                    RegisterOperand dst_rc,
                                     RegisterClass addr_rc,
                                     string dns="">
-  : MIMG_gfx90a <op.GFX10M, (outs 
getLdStRegisterOperandAlign2<dst_rc>.ret:$vdata), dns> {
+  : MIMG_gfx90a <op.GFX10M, (outs getAlign2RegOp<dst_rc>.ret:$vdata), dns> {
   let InOperandList = !con((ins addr_rc:$vaddr, SReg_256_XNULL:$srsrc,
                                 DMask:$dmask, UNorm:$unorm, CPol:$cpol,
                                 R128A16:$r128, LWE:$lwe, DA:$da),
@@ -446,7 +446,7 @@ class MIMG_NoSampler_Helper_gfx90a <mimgopc op, string asm,
 }
 
 class MIMG_NoSampler_gfx10<mimgopc op, string opcode,
-                           RegisterClass DataRC, RegisterClass AddrRC,
+                           RegisterOperand DataRC, RegisterClass AddrRC,
                            string dns="">
   : MIMG_gfx10<op.GFX10M, (outs DataRC:$vdata), dns> {
   let InOperandList = !con((ins AddrRC:$vaddr0, SReg_256_XNULL:$srsrc, 
DMask:$dmask,
@@ -458,7 +458,7 @@ class MIMG_NoSampler_gfx10<mimgopc op, string opcode,
 }
 
 class MIMG_NoSampler_nsa_gfx10<mimgopc op, string opcode,
-                               RegisterClass DataRC, int num_addrs,
+                               RegisterOperand DataRC, int num_addrs,
                                string dns="">
   : MIMG_nsa_gfx10<op.GFX10M, (outs DataRC:$vdata), num_addrs, dns> {
   let InOperandList = !con(AddrIns,
@@ -471,7 +471,7 @@ class MIMG_NoSampler_nsa_gfx10<mimgopc op, string opcode,
 }
 
 class MIMG_NoSampler_gfx11<mimgopc op, string opcode,
-                           RegisterClass DataRC, RegisterClass AddrRC,
+                           RegisterOperand DataRC, RegisterClass AddrRC,
                            string dns="">
   : MIMG_gfx11<op.GFX11, (outs DataRC:$vdata), dns> {
   let InOperandList = !con((ins AddrRC:$vaddr0, SReg_256_XNULL:$srsrc, 
DMask:$dmask,
@@ -483,7 +483,7 @@ class MIMG_NoSampler_gfx11<mimgopc op, string opcode,
 }
 
 class MIMG_NoSampler_nsa_gfx11<mimgopc op, string opcode,
-                               RegisterClass DataRC, int num_addrs,
+                               RegisterOperand DataRC, int num_addrs,
                                string dns="">
   : MIMG_nsa_gfx11<op.GFX11, (outs DataRC:$vdata), num_addrs, dns> {
   let InOperandList = !con(AddrIns,
@@ -496,7 +496,7 @@ class MIMG_NoSampler_nsa_gfx11<mimgopc op, string opcode,
 }
 
 class VIMAGE_NoSampler_gfx12<mimgopc op, string opcode,
-                             RegisterClass DataRC, int num_addrs,
+                             RegisterOperand DataRC, int num_addrs,
                              string dns="">
   : VIMAGE_gfx12<op.GFX12, (outs DataRC:$vdata), num_addrs, dns> {
   let InOperandList = !con(AddrIns,
@@ -507,7 +507,7 @@ class VIMAGE_NoSampler_gfx12<mimgopc op, string opcode,
                     #!if(BaseOpcode.HasD16, "$d16", "");
 }
 
-class VSAMPLE_Sampler_gfx12<mimgopc op, string opcode, RegisterClass DataRC,
+class VSAMPLE_Sampler_gfx12<mimgopc op, string opcode, RegisterOperand DataRC,
                             int num_addrs, RegisterClass Addr3RC = VGPR_32,
                             string dns="">
   : VSAMPLE_gfx12<op.GFX12, (outs DataRC:$vdata), num_addrs, dns, Addr3RC> {
@@ -544,7 +544,7 @@ class VSAMPLE_Sampler_nortn_gfx12<mimgopc op, string opcode,
 }
 
 multiclass MIMG_NoSampler_Src_Helper <mimgopc op, string asm,
-                                      RegisterClass dst_rc, bit enableDisasm,
+                                      RegisterOperand dst_rc, bit enableDisasm,
                                       bit ExtendedImageInst = 1,
                                       bit isVSample = 0> {
   let VAddrDwords = 1 in {
@@ -664,20 +664,20 @@ multiclass MIMG_NoSampler <mimgopc op, string asm, bit 
has_d16, bit mip = 0,
   let BaseOpcode = !cast<MIMGBaseOpcode>(NAME),
       mayLoad = !not(isResInfo) in {
     let VDataDwords = 1 in
-    defm _V1 : MIMG_NoSampler_Src_Helper <op, asm, VGPR_32, 1, msaa>;
+    defm _V1 : MIMG_NoSampler_Src_Helper <op, asm, AVLdSt_32, 1, msaa>;
     let VDataDwords = 2 in
-    defm _V2 : MIMG_NoSampler_Src_Helper <op, asm, VReg_64, 0, msaa>;
+    defm _V2 : MIMG_NoSampler_Src_Helper <op, asm, AVLdSt_64, 0, msaa>;
     let VDataDwords = 3 in
-    defm _V3 : MIMG_NoSampler_Src_Helper <op, asm, VReg_96, 0, msaa>;
+    defm _V3 : MIMG_NoSampler_Src_Helper <op, asm, AVLdSt_96, 0, msaa>;
     let VDataDwords = 4 in
-    defm _V4 : MIMG_NoSampler_Src_Helper <op, asm, VReg_128, 0, msaa>;
+    defm _V4 : MIMG_NoSampler_Src_Helper <op, asm, AVLdSt_128, 0, msaa>;
     let VDataDwords = 5 in
-    defm _V5 : MIMG_NoSampler_Src_Helper <op, asm, VReg_160, 0, msaa>;
+    defm _V5 : MIMG_NoSampler_Src_Helper <op, asm, AVLdSt_160, 0, msaa>;
   }
 }
 
 class MIMG_Store_Helper <mimgopc op, string asm,
-                         RegisterClass data_rc,
+                         RegisterOperand data_rc,
                          RegisterClass addr_rc,
                          string dns = "">
   : MIMG_gfx6789<op.GFX10M, (outs), dns> {
@@ -690,11 +690,11 @@ class MIMG_Store_Helper <mimgopc op, string asm,
 }
 
 class MIMG_Store_Helper_gfx90a <mimgopc op, string asm,
-                                RegisterClass data_rc,
+                                RegisterOperand data_rc,
                                 RegisterClass addr_rc,
                                 string dns = "">
   : MIMG_gfx90a<op.GFX10M, (outs), dns> {
-  let InOperandList = !con((ins 
getLdStRegisterOperandAlign2<data_rc>.ret:$vdata,
+  let InOperandList = !con((ins getAlign2RegOp<data_rc>.ret:$vdata,
                                 addr_rc:$vaddr, SReg_256_XNULL:$srsrc,
                                 DMask:$dmask, UNorm:$unorm, CPol:$cpol,
                                 R128A16:$r128, LWE:$lwe, DA:$da),
@@ -704,7 +704,7 @@ class MIMG_Store_Helper_gfx90a <mimgopc op, string asm,
 }
 
 class MIMG_Store_gfx10<mimgopc op, string opcode,
-                       RegisterClass DataRC, RegisterClass AddrRC,
+                       RegisterOperand DataRC, RegisterClass AddrRC,
                        string dns="">
   : MIMG_gfx10<op.GFX10M, (outs), dns> {
   let InOperandList = !con((ins DataRC:$vdata, AddrRC:$vaddr0, 
SReg_256_XNULL:$srsrc,
@@ -716,7 +716,7 @@ class MIMG_Store_gfx10<mimgopc op, string opcode,
 }
 
 class MIMG_Store_nsa_gfx10<mimgopc op, string opcode,
-                           RegisterClass DataRC, int num_addrs,
+                           RegisterOperand DataRC, int num_addrs,
                            string dns="">
   : MIMG_nsa_gfx10<op.GFX10M, (outs), num_addrs, dns> {
   let InOperandList = !con((ins DataRC:$vdata),
@@ -730,7 +730,7 @@ class MIMG_Store_nsa_gfx10<mimgopc op, string opcode,
 }
 
 class MIMG_Store_gfx11<mimgopc op, string opcode,
-                       RegisterClass DataRC, RegisterClass AddrRC,
+                       RegisterOperand DataRC, RegisterClass AddrRC,
                        string dns="">
   : MIMG_gfx11<op.GFX11, (outs), dns> {
   let InOperandList = !con((ins DataRC:$vdata, AddrRC:$vaddr0, 
SReg_256_XNULL:$srsrc,
@@ -742,7 +742,7 @@ class MIMG_Store_gfx11<mimgopc op, string opcode,
 }
 
 class MIMG_Store_nsa_gfx11<mimgopc op, string opcode,
-                           RegisterClass DataRC, int num_addrs,
+                           RegisterOperand DataRC, int num_addrs,
                            string dns="">
   : MIMG_nsa_gfx11<op.GFX11, (outs), num_addrs, dns> {
   let InOperandList = !con((ins DataRC:$vdata),
@@ -756,7 +756,7 @@ class MIMG_Store_nsa_gfx11<mimgopc op, string opcode,
 }
 
 class VIMAGE_Store_gfx12<mimgopc op, string opcode,
-                         RegisterClass DataRC, int num_addrs,
+                         RegisterOperand DataRC, int num_addrs,
                          string dns="">
   : VIMAGE_gfx12<op.GFX12, (outs), num_addrs, dns> {
   let InOperandList = !con((ins DataRC:$vdata),
@@ -769,7 +769,7 @@ class VIMAGE_Store_gfx12<mimgopc op, string opcode,
 }
 
 multiclass MIMG_Store_Addr_Helper <mimgopc op, string asm,
-                                  RegisterClass data_rc,
+                                  RegisterOperand data_rc,
                                   bit enableDisasm> {
   let mayLoad = 0, mayStore = 1, hasSideEffects = 0, hasPostISelHook = 0,
       DisableWQM = 1 in {
@@ -860,19 +860,19 @@ multiclass MIMG_Store <mimgopc op, string asm, bit 
has_d16, bit mip = 0> {
 
   let BaseOpcode = !cast<MIMGBaseOpcode>(NAME) in {
     let VDataDwords = 1 in
-    defm _V1 : MIMG_Store_Addr_Helper <op, asm, VGPR_32, 1>;
+    defm _V1 : MIMG_Store_Addr_Helper <op, asm, AVLdSt_32, 1>;
     let VDataDwords = 2 in
-    defm _V2 : MIMG_Store_Addr_Helper <op, asm, VReg_64, 0>;
+    defm _V2 : MIMG_Store_Addr_Helper <op, asm, AVLdSt_64, 0>;
     let VDataDwords = 3 in
-    defm _V3 : MIMG_Store_Addr_Helper <op, asm, VReg_96, 0>;
+    defm _V3 : MIMG_Store_Addr_Helper <op, asm, AVLdSt_96, 0>;
     let VDataDwords = 4 in
-    defm _V4 : MIMG_Store_Addr_Helper <op, asm, VReg_128, 0>;
+    defm _V4 : MIMG_Store_Addr_Helper <op, asm, AVLdSt_128, 0>;
     let VDataDwords = 5 in
-    defm _V5 : MIMG_Store_Addr_Helper <op, asm, VReg_160, 0>;
+    defm _V5 : MIMG_Store_Addr_Helper <op, asm, AVLdSt_160, 0>;
   }
 }
 
-class MIMG_Atomic_gfx6789_base <bits<8> op, string asm, RegisterClass data_rc,
+class MIMG_Atomic_gfx6789_base <bits<8> op, string asm, RegisterOperand 
data_rc,
                                 RegisterClass addr_rc, string dns="">
   : MIMG_gfx6789 <op, (outs data_rc:$vdst), dns> {
   let Constraints = "$vdst = $vdata";
@@ -883,33 +883,33 @@ class MIMG_Atomic_gfx6789_base <bits<8> op, string asm, 
RegisterClass data_rc,
   let AsmString = asm#" $vdst, $vaddr, 
$srsrc$dmask$unorm$cpol$r128$tfe$lwe$da";
 }
 
-class MIMG_Atomic_gfx90a_base <bits<8> op, string asm, RegisterClass data_rc,
+class MIMG_Atomic_gfx90a_base <bits<8> op, string asm, RegisterOperand data_rc,
                                RegisterClass addr_rc, string dns="">
-  : MIMG_gfx90a <op, (outs getLdStRegisterOperandAlign2<data_rc>.ret:$vdst), 
dns> {
+  : MIMG_gfx90a <op, (outs getAlign2RegOp<data_rc>.ret:$vdst), dns> {
   let Constraints = "$vdst = $vdata";
 
-  let InOperandList = (ins getLdStRegisterOperandAlign2<data_rc>.ret:$vdata,
+  let InOperandList = (ins getAlign2RegOp<data_rc>.ret:$vdata,
                            addr_rc:$vaddr, SReg_256_XNULL:$srsrc,
                            DMask:$dmask, UNorm:$unorm, CPol:$cpol,
                            R128A16:$r128, LWE:$lwe, DA:$da);
   let AsmString = asm#" $vdst, $vaddr, $srsrc$dmask$unorm$cpol$r128$lwe$da";
 }
 
-class MIMG_Atomic_si<mimgopc op, string asm, RegisterClass data_rc,
+class MIMG_Atomic_si<mimgopc op, string asm, RegisterOperand data_rc,
                      RegisterClass addr_rc, bit enableDasm = 0>
   : MIMG_Atomic_gfx6789_base<op.SI, asm, data_rc, addr_rc,
                              !if(enableDasm, "GFX6GFX7", "")> {
   let AssemblerPredicate = isGFX6GFX7;
 }
 
-class MIMG_Atomic_vi<mimgopc op, string asm, RegisterClass data_rc,
+class MIMG_Atomic_vi<mimgopc op, string asm, RegisterOperand data_rc,
                      RegisterClass addr_rc, bit enableDasm = 0>
   : MIMG_Atomic_gfx6789_base<op.VI, asm, data_rc, addr_rc, !if(enableDasm, 
"GFX8", "")> {
   let AssemblerPredicate = isGFX8GFX9NotGFX90A;
   let MIMGEncoding = MIMGEncGfx8;
 }
 
-class MIMG_Atomic_gfx90a<mimgopc op, string asm, RegisterClass data_rc,
+class MIMG_Atomic_gfx90a<mimgopc op, string asm, RegisterOperand data_rc,
                          RegisterClass addr_rc, bit enableDasm = 0>
   : MIMG_Atomic_gfx90a_base<op.VI, asm, data_rc, addr_rc, !if(enableDasm, 
"GFX90A", "")> {
   let AssemblerPredicate = isGFX90APlus;
@@ -917,7 +917,7 @@ class MIMG_Atomic_gfx90a<mimgopc op, string asm, 
RegisterClass data_rc,
 }
 
 class MIMG_Atomic_gfx10<mimgopc op, string opcode,
-                        RegisterClass DataRC, RegisterClass AddrRC,
+                        RegisterOperand DataRC, RegisterClass AddrRC,
                         bit enableDisasm = 0>
   : MIMG_gfx10<op.GFX10M, (outs DataRC:$vdst),
                !if(enableDisasm, "GFX10", "")> {
@@ -930,7 +930,7 @@ class MIMG_Atomic_gfx10<mimgopc op, string opcode,
 }
 
 class MIMG_Atomic_nsa_gfx10<mimgopc op, string opcode,
-                            RegisterClass DataRC, int num_addrs,
+                            RegisterOperand DataRC, int num_addrs,
                             bit enableDisasm = 0>
   : MIMG_nsa_gfx10<op.GFX10M, (outs DataRC:$vdst), num_addrs,
                    !if(enableDisasm, "GFX10", "")> {
@@ -945,7 +945,7 @@ class MIMG_Atomic_nsa_gfx10<mimgopc op, string opcode,
 }
 
 class MIMG_Atomic_gfx11<mimgopc op, string opcode,
-                        RegisterClass DataRC, RegisterClass AddrRC,
+                        RegisterOperand DataRC, RegisterClass AddrRC,
                         bit enableDisasm = 0>
   : MIMG_gfx11<op.GFX11, (outs DataRC:$vdst),
                !if(enableDisasm, "GFX11", "")> {
@@ -958,7 +958,7 @@ class MIMG_Atomic_gfx11<mimgopc op, string opcode,
 }
 
 class MIMG_Atomic_nsa_gfx11<mimgopc op, string opcode,
-                            RegisterClass DataRC, int num_addrs,
+                            RegisterOperand DataRC, int num_addrs,
                             bit enableDisasm = 0>
   : MIMG_nsa_gfx11<op.GFX11, (outs DataRC:$vdst), num_addrs,
                    !if(enableDisasm, "GFX11", "")> {
@@ -972,7 +972,7 @@ class MIMG_Atomic_nsa_gfx11<mimgopc op, string opcode,
   let AsmString = opcode#" $vdata, "#AddrAsm#", 
$srsrc$dmask$dim$unorm$cpol$r128$a16$tfe$lwe";
 }
 
-class VIMAGE_Atomic_gfx12<mimgopc op, string opcode, RegisterClass DataRC,
+class VIMAGE_Atomic_gfx12<mimgopc op, string opcode, RegisterOperand DataRC,
                           int num_addrs, string renamed, bit enableDisasm = 0>
   : VIMAGE_gfx12<op.GFX12, (outs DataRC:$vdst), num_addrs,
                   !if(enableDisasm, "GFX12", "")> {
@@ -987,7 +987,7 @@ class VIMAGE_Atomic_gfx12<mimgopc op, string opcode, 
RegisterClass DataRC,
 }
 
 multiclass MIMG_Atomic_Addr_Helper_m <mimgopc op, string asm,
-                                      RegisterClass data_rc,
+                                      RegisterOperand data_rc,
                                       bit enableDasm = 0,
                                       bit isFP = 0,
                                       string renamed = ""> {
@@ -1105,19 +1105,19 @@ multiclass MIMG_Atomic <mimgopc op, string asm, bit 
isCmpSwap = 0, bit isFP = 0,
       // Other variants are reconstructed by disassembler using dmask and tfe.
       if !not(isCmpSwap) then {
         let VDataDwords = 1 in
-        defm _V1 : MIMG_Atomic_Addr_Helper_m <op, asm, VGPR_32, 1, isFP, 
renamed>;
+        defm _V1 : MIMG_Atomic_Addr_Helper_m <op, asm, AVLdSt_32, 1, isFP, 
renamed>;
       }
 
       let VDataDwords = 2 in
-      defm _V2 : MIMG_Atomic_Addr_Helper_m <op, asm, VReg_64, isCmpSwap, isFP, 
renamed>;
+      defm _V2 : MIMG_Atomic_Addr_Helper_m <op, asm, AVLdSt_64, isCmpSwap, 
isFP, renamed>;
       let VDataDwords = 3 in
-      defm _V3 : MIMG_Atomic_Addr_Helper_m <op, asm, VReg_96, 0, isFP, 
renamed>;
+      defm _V3 : MIMG_Atomic_Addr_Helper_m <op, asm, AVLdSt_96, 0, isFP, 
renamed>;
 
       if isCmpSwap then {
         let VDataDwords = 4 in
-        defm _V4 : MIMG_Atomic_Addr_Helper_m <op, asm, VReg_128, 0, isFP, 
renamed>;
+        defm _V4 : MIMG_Atomic_Addr_Helper_m <op, asm, AVLdSt_128, 0, isFP, 
renamed>;
         let VDataDwords = 5 in
-        defm _V5 : MIMG_Atomic_Addr_Helper_m <op, asm, VReg_160, 0, isFP, 
renamed>;
+        defm _V5 : MIMG_Atomic_Addr_Helper_m <op, asm, AVLdSt_160, 0, isFP, 
renamed>;
       }
     }
   } // End IsAtomicRet = 1
@@ -1127,7 +1127,7 @@ multiclass MIMG_Atomic_Renamed <mimgopc op, string asm, 
string renamed,
                                 bit isCmpSwap = 0, bit isFP = 0>
   : MIMG_Atomic <op, asm, isCmpSwap, isFP, renamed>;
 
-class MIMG_Sampler_Helper <mimgopc op, string asm, RegisterClass dst_rc,
+class MIMG_Sampler_Helper <mimgopc op, string asm, RegisterOperand dst_rc,
                            RegisterClass src_rc, string dns="">
   : MIMG_gfx6789 <op.VI, (outs dst_rc:$vdata), dns> {
   let InOperandList = !con((ins src_rc:$vaddr, SReg_256_XNULL:$srsrc, 
SReg_128_XNULL:$ssamp,
@@ -1138,9 +1138,9 @@ class MIMG_Sampler_Helper <mimgopc op, string asm, 
RegisterClass dst_rc,
                       #!if(BaseOpcode.HasD16, "$d16", "");
 }
 
-class MIMG_Sampler_gfx90a<mimgopc op, string asm, RegisterClass dst_rc,
+class MIMG_Sampler_gfx90a<mimgopc op, string asm, RegisterOperand dst_rc,
                           RegisterClass src_rc, string dns="">
-  : MIMG_gfx90a<op.GFX10M, (outs 
getLdStRegisterOperandAlign2<dst_rc>.ret:$vdata), dns> {
+  : MIMG_gfx90a<op.GFX10M, (outs dst_rc:$vdata), dns> {
   let InOperandList = !con((ins src_rc:$vaddr, SReg_256_XNULL:$srsrc, 
SReg_128_XNULL:$ssamp,
                                 DMask:$dmask, UNorm:$unorm, CPol:$cpol,
                                 R128A16:$r128, LWE:$lwe, DA:$da),
@@ -1164,7 +1164,7 @@ class MIMG_Sampler_Asm_gfx10p<string opcode, string 
AsmPrefix, bit HasD16> {
 }
 
 class MIMG_Sampler_gfx10<mimgopc op, string opcode,
-                         RegisterClass DataRC, RegisterClass AddrRC,
+                         RegisterOperand DataRC, RegisterClass AddrRC,
                          string dns="">
   : MIMG_gfx10<op.GFX10M, (outs DataRC:$vdata), dns> {
   let InOperandList = MIMG_Sampler_OpList_gfx10p<(ins AddrRC:$vaddr0), 
BaseOpcode.HasD16>.ret;
@@ -1172,7 +1172,7 @@ class MIMG_Sampler_gfx10<mimgopc op, string opcode,
 }
 
 class MIMG_Sampler_nsa_gfx10<mimgopc op, string opcode,
-                             RegisterClass DataRC, int num_addrs,
+                             RegisterOperand DataRC, int num_addrs,
                              string dns="">
   : MIMG_nsa_gfx10<op.GFX10M, (outs DataRC:$vdata), num_addrs, dns> {
   let InOperandList = MIMG_Sampler_OpList_gfx10p<AddrIns, 
BaseOpcode.HasD16>.ret;
@@ -1200,7 +1200,7 @@ class MIMG_Sampler_nortn_nsa_gfx10<mimgopc op, string 
opcode,
 }
 
 class MIMG_Sampler_gfx11<mimgopc op, string opcode,
-                         RegisterClass DataRC, RegisterClass AddrRC,
+                         RegisterOperand DataRC, RegisterClass AddrRC,
                          string dns="">
   : MIMG_gfx11<op.GFX11, (outs DataRC:$vdata), dns> {
   let InOperandList = MIMG_Sampler_OpList_gfx10p<(ins AddrRC:$vaddr0), 
BaseOpcode.HasD16>.ret;
@@ -1208,7 +1208,7 @@ class MIMG_Sampler_gfx11<mimgopc op, string opcode,
 }
 
 class MIMG_Sampler_nsa_gfx11<mimgopc op, string opcode,
-                             RegisterClass DataRC, int num_addrs,
+                             RegisterOperand DataRC, int num_addrs,
                              RegisterClass LastVAddrSize, string dns="">
   : MIMG_nsa_gfx11<op.GFX11, (outs DataRC:$vdata), num_addrs, dns, [],
                    LastVAddrSize> {
@@ -1345,7 +1345,7 @@ class MIMG_Sampler_AddrSizes<AMDGPUSampleVariant sample, 
bit isG16,
 }
 
 multiclass MIMG_Sampler_Src_Helper <mimgopc op, string asm,
-                                    AMDGPUSampleVariant sample, RegisterClass 
dst_rc,
+                                    AMDGPUSampleVariant sample, 
RegisterOperand dst_rc,
                                     bit enableDisasm = 0,
                                     bit ExtendedImageInst = 1, bit isG16 = 0> {
   foreach addr = MIMG_Sampler_AddrSizes<sample, isG16>.MachineInstrs in {
@@ -1473,15 +1473,15 @@ multiclass MIMG_Sampler <mimgopc op, 
AMDGPUSampleVariant sample, bit isPointSamp
   let BaseOpcode = !cast<MIMGBaseOpcode>(NAME), WQM = wqm,
       mayLoad = !not(isGetLod) in {
     let VDataDwords = 1 in
-    defm _V1 : MIMG_Sampler_Src_Helper<op, asm, sample, VGPR_32, 1, 
ExtendedImageInst, isG16>;
+    defm _V1 : MIMG_Sampler_Src_Helper<op, asm, sample, AVLdSt_32, 1, 
ExtendedImageInst, isG16>;
     let VDataDwords = 2 in
-    defm _V2 : MIMG_Sampler_Src_Helper<op, asm, sample, VReg_64, 0, 
ExtendedImageInst, isG16>;
+    defm _V2 : MIMG_Sampler_Src_Helper<op, asm, sample, AVLdSt_64, 0, 
ExtendedImageInst, isG16>;
     let VDataDwords = 3 in
-    defm _V3 : MIMG_Sampler_Src_Helper<op, asm, sample, VReg_96, 0, 
ExtendedImageInst, isG16>;
+    defm _V3 : MIMG_Sampler_Src_Helper<op, asm, sample, AVLdSt_96, 0, 
ExtendedImageInst, isG16>;
     let VDataDwords = 4 in
-    defm _V4 : MIMG_Sampler_Src_Helper<op, asm, sample, VReg_128, 0, 
ExtendedImageInst, isG16>;
+    defm _V4 : MIMG_Sampler_Src_Helper<op, asm, sample, AVLdSt_128, 0, 
ExtendedImageInst, isG16>;
     let VDataDwords = 5 in
-    defm _V5 : MIMG_Sampler_Src_Helper<op, asm, sample, VReg_160, 0, 
ExtendedImageInst, isG16>;
+    defm _V5 : MIMG_Sampler_Src_Helper<op, asm, sample, AVLdSt_160, 0, 
ExtendedImageInst, isG16>;
   }
 
   if !not(isGetLod) then
@@ -1501,11 +1501,11 @@ multiclass MIMG_Gather <mimgopc op, AMDGPUSampleVariant 
sample, bit wqm = 0,
   let BaseOpcode = !cast<MIMGBaseOpcode>(NAME), WQM = wqm,
       Gather4 = 1 in {
     let VDataDwords = 2 in
-    defm _V2 : MIMG_Sampler_Src_Helper<op, asm, sample, VReg_64, 
/*enableDisasm*/ true>; /* for packed D16 only */
+    defm _V2 : MIMG_Sampler_Src_Helper<op, asm, sample, AVLdSt_64, 
/*enableDisasm*/ true>; /* for packed D16 only */
     let VDataDwords = 4 in
-    defm _V4 : MIMG_Sampler_Src_Helper<op, asm, sample, VReg_128>;
+    defm _V4 : MIMG_Sampler_Src_Helper<op, asm, sample, AVLdSt_128>;
     let VDataDwords = 5 in
-    defm _V5 : MIMG_Sampler_Src_Helper<op, asm, sample, VReg_160>;
+    defm _V5 : MIMG_Sampler_Src_Helper<op, asm, sample, AVLdSt_160>;
   }
 }
 
@@ -1632,13 +1632,13 @@ multiclass MIMG_MSAA_Load <mimgopc op, string asm> {
   let BaseOpcode = !cast<MIMGBaseOpcode>(NAME),
     Gather4 = 1, hasPostISelHook = 0, mayLoad = 1 in {
     let VDataDwords = 2 in
-    defm _V2 : MIMG_NoSampler_Src_Helper<op, asm, VReg_64, 0, 0, 1>; /* packed 
D16 */
+    defm _V2 : MIMG_NoSampler_Src_Helper<op, asm, AVLdSt_64, 0, 0, 1>; /* 
packed D16 */
     let VDataDwords = 3 in
-    defm _V3 : MIMG_NoSampler_Src_Helper<op, asm, VReg_96, 0, 0, 1>; /* packed 
D16 + tfe */
+    defm _V3 : MIMG_NoSampler_Src_Helper<op, asm, AVLdSt_96, 0, 0, 1>; /* 
packed D16 + tfe */
     let VDataDwords = 4 in
-    defm _V4 : MIMG_NoSampler_Src_Helper<op, asm, VReg_128, 1, 0, 1>;
+    defm _V4 : MIMG_NoSampler_Src_Helper<op, asm, AVLdSt_128, 1, 0, 1>;
     let VDataDwords = 5 in
-    defm _V5 : MIMG_NoSampler_Src_Helper<op, asm, VReg_160, 0, 0, 1>;
+    defm _V5 : MIMG_NoSampler_Src_Helper<op, asm, AVLdSt_160, 0, 0, 1>;
   }
 }
 
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.td 
b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
index 6e068fbd6a10e..72b602b7c42c6 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.td
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
@@ -2604,21 +2604,28 @@ class getLdStRegisterOperandForVT<ValueType VT> {
   RegisterOperand ret = getLdStRegisterOperandForSize<VT.Size>.ret;
 }
 
-class getLdStRegisterOperandAlign2<RegisterClass RC> {
-  // This type of operands is only used in pseudo instructions helping
-  // code generation and thus doesn't need encoding and decoding methods.
-  // It also doesn't need to support AGPRs, because GFX908/A/40 do not
-  // support True16.
-  defvar VLdSt_16 = RegisterOperand<VGPR_16>;
-
+class getAlign2RegOp<RegisterOperand RC> {
   RegisterOperand ret =
-    !cond(!eq(RC.Size, 16)   : VLdSt_16,
-          !eq(RC.Size, 32)   : AVLdSt_32,
-          !eq(RC.Size, 64)   : AVLdSt_64_Align2,
-          !eq(RC.Size, 96)   : AVLdSt_96_Align2,
-          !eq(RC.Size, 128)  : AVLdSt_128_Align2,
-          !eq(RC.Size, 160)  : AVLdSt_160_Align2,
-          !eq(RC.Size, 1024) : AVLdSt_1024_Align2);
+    !cond(!eq(RC, VGPROp_16) : VGPROp_16,
+          !eq(RC, VGPROp_32) : VGPROp_32,
+          !eq(RC, VGPROp_64) : VGPROp_64_Align2,
+          !eq(RC, VGPROp_64_Align1) : VGPROp_64_Align2,
+          !eq(RC, VGPROp_96) : VGPROp_96_Align2,
+          !eq(RC, VGPROp_96_Align1) : VGPROp_96_Align2,
+          !eq(RC, VGPROp_128) : VGPROp_128_Align2,
+          !eq(RC, VGPROp_128_Align1) : VGPROp_128_Align2,
+          !eq(RC, VGPROp_160) : VGPROp_160_Align2,
+          !eq(RC, VGPROp_160_Align1) : VGPROp_160_Align2,
+          !eq(RC, VGPROp_1024) : VGPROp_1024_Align2,
+          !eq(RC, VGPROp_1024_Align1) : VGPROp_1024_Align2,
+          !eq(RC, AVLdSt_32) : AVLdSt_32,
+          !eq(RC, AVLdSt_64) : AVLdSt_64_Align2,
+          !eq(RC, AVLdSt_96) : AVLdSt_96_Align2,
+          !eq(RC, AVLdSt_96_Align1) : AVLdSt_96_Align2,
+          !eq(RC, AVLdSt_128) : AVLdSt_128_Align2,
+          !eq(RC, AVLdSt_128_Align1) : AVLdSt_128_Align2,
+          !eq(RC, AVLdSt_160) : AVLdSt_160_Align2,
+          !eq(RC, AVLdSt_160_Align1) : AVLdSt_160_Align2);
 }
 
 class getEquivalentAGPRClass<RegisterClass RC> {
diff --git a/llvm/lib/Target/AMDGPU/SIRegisterInfo.td 
b/llvm/lib/Target/AMDGPU/SIRegisterInfo.td
index a8e330acbc81e..162ef647be7e9 100644
--- a/llvm/lib/Target/AMDGPU/SIRegisterInfo.td
+++ b/llvm/lib/Target/AMDGPU/SIRegisterInfo.td
@@ -1348,7 +1348,7 @@ def VGPRSrc_16 : RegisterOperand<VGPR_16> {
 def VGPROp_16 : RegisterOperand<VGPR_16>;
 def VGPROp_32 : RegisterOperand<VGPR_32>;
 
-foreach size = ["64", "96", "128", "256", "1024" ] in {
+foreach size = ["64", "96", "128", "160", "256", "1024" ] in {
   def VGPROp_#size : RegisterOperand<!cast<RegisterClass>("VReg_"#size)>;
   def VGPROp_#size#_Align1 : 
RegisterOperand<!cast<RegisterClass>("VReg_"#size)>;
   def VGPROp_#size#_Align2 : 
RegisterOperand<!cast<RegisterClass>("VReg_"#size#_Align2)>;
diff --git a/llvm/test/MC/AMDGPU/gfx90a_ldst_acc.s 
b/llvm/test/MC/AMDGPU/gfx90a_ldst_acc.s
index 002ae4bc88e1b..43673d1d49c79 100644
--- a/llvm/test/MC/AMDGPU/gfx90a_ldst_acc.s
+++ b/llvm/test/MC/AMDGPU/gfx90a_ldst_acc.s
@@ -9978,725 +9978,725 @@ ds_read_b128 a[6:9], v1
 ds_read_b128 a[6:9], v1 offset:4
 
 // GFX90A: image_load a5, v[2:5], s[8:15] dmask:0x1 ; encoding: 
[0x00,0x01,0x01,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a5, v[2:5], s[8:15] dmask:0x1
 
 // GFX90A: image_load a252, v[2:5], s[8:15] dmask:0x1 ; encoding: 
[0x00,0x01,0x01,0xf0,0x02,0xfc,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a252, v[2:5], s[8:15] dmask:0x1
 
 // GFX90A: image_load a5, v[252:255], s[8:15] dmask:0x1 ; encoding: 
[0x00,0x01,0x01,0xf0,0xfc,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a5, v[252:255], s[8:15] dmask:0x1
 
 // GFX90A: image_load a5, v[2:5], s[12:19] dmask:0x1 ; encoding: 
[0x00,0x01,0x01,0xf0,0x02,0x05,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a5, v[2:5], s[12:19] dmask:0x1
 
 // GFX90A: image_load a5, v[2:5], s[92:99] dmask:0x1 ; encoding: 
[0x00,0x01,0x01,0xf0,0x02,0x05,0x17,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a5, v[2:5], s[92:99] dmask:0x1
 
 // GFX90A: image_load a5, v[2:5], s[8:15] dmask:0x2 ; encoding: 
[0x00,0x02,0x01,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a5, v[2:5], s[8:15] dmask:0x2
 
 // GFX90A: image_load a[6:7], v[2:5], s[8:15] dmask:0x3 ; encoding: 
[0x00,0x03,0x01,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a[6:7], v[2:5], s[8:15] dmask:0x3
 
 // GFX90A: image_load a5, v[2:5], s[8:15] dmask:0x4 ; encoding: 
[0x00,0x04,0x01,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a5, v[2:5], s[8:15] dmask:0x4
 
 // GFX90A: image_load a[6:7], v[2:5], s[8:15] dmask:0x5 ; encoding: 
[0x00,0x05,0x01,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a[6:7], v[2:5], s[8:15] dmask:0x5
 
 // GFX90A: image_load a[6:7], v[2:5], s[8:15] dmask:0x6 ; encoding: 
[0x00,0x06,0x01,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a[6:7], v[2:5], s[8:15] dmask:0x6
 
 // GFX90A: image_load a[6:8], v[2:5], s[8:15] dmask:0x7 ; encoding: 
[0x00,0x07,0x01,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a[6:8], v[2:5], s[8:15] dmask:0x7
 
 // GFX90A: image_load a5, v[2:5], s[8:15] dmask:0x8 ; encoding: 
[0x00,0x08,0x01,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a5, v[2:5], s[8:15] dmask:0x8
 
 // GFX90A: image_load a[6:7], v[2:5], s[8:15] dmask:0x9 ; encoding: 
[0x00,0x09,0x01,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a[6:7], v[2:5], s[8:15] dmask:0x9
 
 // GFX90A: image_load a[6:7], v[2:5], s[8:15] dmask:0xa ; encoding: 
[0x00,0x0a,0x01,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a[6:7], v[2:5], s[8:15] dmask:0xa
 
 // GFX90A: image_load a[6:8], v[2:5], s[8:15] dmask:0xb ; encoding: 
[0x00,0x0b,0x01,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a[6:8], v[2:5], s[8:15] dmask:0xb
 
 // GFX90A: image_load a[6:7], v[2:5], s[8:15] dmask:0xc ; encoding: 
[0x00,0x0c,0x01,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a[6:7], v[2:5], s[8:15] dmask:0xc
 
 // GFX90A: image_load a[6:8], v[2:5], s[8:15] dmask:0xd ; encoding: 
[0x00,0x0d,0x01,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a[6:8], v[2:5], s[8:15] dmask:0xd
 
 // GFX90A: image_load a[6:8], v[2:5], s[8:15] dmask:0xe ; encoding: 
[0x00,0x0e,0x01,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a[6:8], v[2:5], s[8:15] dmask:0xe
 
 // GFX90A: image_load a5, v[2:5], s[8:15]  ; encoding: 
[0x00,0x00,0x01,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a5, v[2:5], s[8:15]
 
 // GFX90A: image_load a5, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x01,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a5, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_load a5, v[2:5], s[8:15] dmask:0x1 glc ; encoding: 
[0x00,0x21,0x01,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a5, v[2:5], s[8:15] dmask:0x1 glc
 
 // GFX90A: image_load a5, v[2:5], s[8:15] dmask:0x1 slc ; encoding: 
[0x00,0x01,0x01,0xf2,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a5, v[2:5], s[8:15] dmask:0x1 slc
 
 // GFX90A: image_load a5, v[2:5], s[8:15] dmask:0x1 lwe ; encoding: 
[0x00,0x01,0x03,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a5, v[2:5], s[8:15] dmask:0x1 lwe
 
 // GFX90A: image_load a5, v[2:5], s[8:15] dmask:0x1 da ; encoding: 
[0x00,0x41,0x01,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a5, v[2:5], s[8:15] dmask:0x1 da
 
 // GFX90A: image_load a5, v[2:5], s[8:15] dmask:0x1 d16 ; encoding: 
[0x00,0x01,0x01,0xf0,0x02,0x05,0x02,0x80]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_load a5, v[2:5], s[8:15] dmask:0x1 d16
 
 // GFX90A: image_store a1, v[2:5], s[12:19] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x21,0xf0,0x02,0x01,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a1, v[2:5], s[12:19] dmask:0x1 unorm
 
 // GFX90A: image_store a252, v[2:5], s[12:19] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x21,0xf0,0x02,0xfc,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a252, v[2:5], s[12:19] dmask:0x1 unorm
 
 // GFX90A: image_store a1, v[252:255], s[12:19] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x21,0xf0,0xfc,0x01,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a1, v[252:255], s[12:19] dmask:0x1 unorm
 
 // GFX90A: image_store a1, v[2:5], s[16:23] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x21,0xf0,0x02,0x01,0x04,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a1, v[2:5], s[16:23] dmask:0x1 unorm
 
 // GFX90A: image_store a1, v[2:5], s[92:99] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x21,0xf0,0x02,0x01,0x17,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a1, v[2:5], s[92:99] dmask:0x1 unorm
 
 // GFX90A: image_store a1, v[2:5], s[12:19] dmask:0x2 unorm ; encoding: 
[0x00,0x12,0x21,0xf0,0x02,0x01,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a1, v[2:5], s[12:19] dmask:0x2 unorm
 
 // GFX90A: image_store a[2:3], v[2:5], s[12:19] dmask:0x3 unorm ; encoding: 
[0x00,0x13,0x21,0xf0,0x02,0x02,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a[2:3], v[2:5], s[12:19] dmask:0x3 unorm
 
 // GFX90A: image_store a1, v[2:5], s[12:19] dmask:0x4 unorm ; encoding: 
[0x00,0x14,0x21,0xf0,0x02,0x01,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a1, v[2:5], s[12:19] dmask:0x4 unorm
 
 // GFX90A: image_store a[2:3], v[2:5], s[12:19] dmask:0x5 unorm ; encoding: 
[0x00,0x15,0x21,0xf0,0x02,0x02,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a[2:3], v[2:5], s[12:19] dmask:0x5 unorm
 
 // GFX90A: image_store a[2:3], v[2:5], s[12:19] dmask:0x6 unorm ; encoding: 
[0x00,0x16,0x21,0xf0,0x02,0x02,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a[2:3], v[2:5], s[12:19] dmask:0x6 unorm
 
 // GFX90A: image_store a[2:4], v[2:5], s[12:19] dmask:0x7 unorm ; encoding: 
[0x00,0x17,0x21,0xf0,0x02,0x02,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a[2:4], v[2:5], s[12:19] dmask:0x7 unorm
 
 // GFX90A: image_store a1, v[2:5], s[12:19] dmask:0x8 unorm ; encoding: 
[0x00,0x18,0x21,0xf0,0x02,0x01,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a1, v[2:5], s[12:19] dmask:0x8 unorm
 
 // GFX90A: image_store a[2:3], v[2:5], s[12:19] dmask:0x9 unorm ; encoding: 
[0x00,0x19,0x21,0xf0,0x02,0x02,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a[2:3], v[2:5], s[12:19] dmask:0x9 unorm
 
 // GFX90A: image_store a[2:3], v[2:5], s[12:19] dmask:0xa unorm ; encoding: 
[0x00,0x1a,0x21,0xf0,0x02,0x02,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a[2:3], v[2:5], s[12:19] dmask:0xa unorm
 
 // GFX90A: image_store a[2:4], v[2:5], s[12:19] dmask:0xb unorm ; encoding: 
[0x00,0x1b,0x21,0xf0,0x02,0x02,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a[2:4], v[2:5], s[12:19] dmask:0xb unorm
 
 // GFX90A: image_store a[2:3], v[2:5], s[12:19] dmask:0xc unorm ; encoding: 
[0x00,0x1c,0x21,0xf0,0x02,0x02,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a[2:3], v[2:5], s[12:19] dmask:0xc unorm
 
 // GFX90A: image_store a[2:4], v[2:5], s[12:19] dmask:0xd unorm ; encoding: 
[0x00,0x1d,0x21,0xf0,0x02,0x02,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a[2:4], v[2:5], s[12:19] dmask:0xd unorm
 
 // GFX90A: image_store a[2:4], v[2:5], s[12:19] dmask:0xe unorm ; encoding: 
[0x00,0x1e,0x21,0xf0,0x02,0x02,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a[2:4], v[2:5], s[12:19] dmask:0xe unorm
 
 // GFX90A: image_store a[2:5], v[2:5], s[12:19] dmask:0xf unorm ; encoding: 
[0x00,0x1f,0x21,0xf0,0x02,0x02,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a[2:5], v[2:5], s[12:19] dmask:0xf unorm
 
 // GFX90A: image_store a1, v[2:5], s[12:19] unorm ; encoding: 
[0x00,0x10,0x21,0xf0,0x02,0x01,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a1, v[2:5], s[12:19] unorm
 
 // GFX90A: image_store a1, v[2:5], s[12:19] dmask:0x1 unorm glc ; encoding: 
[0x00,0x31,0x21,0xf0,0x02,0x01,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a1, v[2:5], s[12:19] dmask:0x1 unorm glc
 
 // GFX90A: image_store a1, v[2:5], s[12:19] dmask:0x1 unorm slc ; encoding: 
[0x00,0x11,0x21,0xf2,0x02,0x01,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a1, v[2:5], s[12:19] dmask:0x1 unorm slc
 
 // GFX90A: image_store a1, v[2:5], s[12:19] dmask:0x1 unorm lwe ; encoding: 
[0x00,0x11,0x23,0xf0,0x02,0x01,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a1, v[2:5], s[12:19] dmask:0x1 unorm lwe
 
 // GFX90A: image_store a1, v[2:5], s[12:19] dmask:0x1 unorm da ; encoding: 
[0x00,0x51,0x21,0xf0,0x02,0x01,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a1, v[2:5], s[12:19] dmask:0x1 unorm da
 
 // GFX90A: image_store a1, v[2:5], s[12:19] dmask:0x1 unorm d16 ; encoding: 
[0x00,0x11,0x21,0xf0,0x02,0x01,0x03,0x80]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_store a1, v[2:5], s[12:19] dmask:0x1 unorm d16
 
 // GFX90A: image_atomic_swap a5, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x41,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_swap a5, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_swap a252, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x41,0xf0,0x02,0xfc,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_swap a252, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_swap a5, v[252:255], s[8:15] dmask:0x1 unorm ; 
encoding: [0x00,0x11,0x41,0xf0,0xfc,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_swap a5, v[252:255], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_swap a5, v[2:5], s[12:19] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x41,0xf0,0x02,0x05,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_swap a5, v[2:5], s[12:19] dmask:0x1 unorm
 
 // GFX90A: image_atomic_swap a5, v[2:5], s[92:99] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x41,0xf0,0x02,0x05,0x17,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_swap a5, v[2:5], s[92:99] dmask:0x1 unorm
 
 // GFX90A: image_atomic_swap a[6:7], v[2:5], s[8:15] dmask:0x3 unorm ; 
encoding: [0x00,0x13,0x41,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_swap a[6:7], v[2:5], s[8:15] dmask:0x3 unorm
 
 // GFX90A: image_atomic_swap a5, v[2:5], s[8:15] dmask:0x1 unorm glc ; 
encoding: [0x00,0x31,0x41,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_swap a5, v[2:5], s[8:15] dmask:0x1 unorm glc
 
 // GFX90A: image_atomic_swap a5, v[2:5], s[8:15] dmask:0x1 unorm slc ; 
encoding: [0x00,0x11,0x41,0xf2,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_swap a5, v[2:5], s[8:15] dmask:0x1 unorm slc
 
 // GFX90A: image_atomic_swap a5, v[2:5], s[8:15] dmask:0x1 unorm lwe ; 
encoding: [0x00,0x11,0x43,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_swap a5, v[2:5], s[8:15] dmask:0x1 unorm lwe
 
 // GFX90A: image_atomic_swap a5, v[2:5], s[8:15] dmask:0x1 unorm da ; 
encoding: [0x00,0x51,0x41,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_swap a5, v[2:5], s[8:15] dmask:0x1 unorm da
 
 // GFX90A: image_atomic_cmpswap a[6:7], v[2:5], s[8:15] dmask:0x3 unorm ; 
encoding: [0x00,0x13,0x45,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_cmpswap a[6:7], v[2:5], s[8:15] dmask:0x3 unorm
 
 // GFX90A: image_atomic_cmpswap a[252:253], v[2:5], s[8:15] dmask:0x3 unorm ; 
encoding: [0x00,0x13,0x45,0xf0,0x02,0xfc,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_cmpswap a[252:253], v[2:5], s[8:15] dmask:0x3 unorm
 
 // GFX90A: image_atomic_cmpswap a[6:7], v[252:255], s[8:15] dmask:0x3 unorm ; 
encoding: [0x00,0x13,0x45,0xf0,0xfc,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_cmpswap a[6:7], v[252:255], s[8:15] dmask:0x3 unorm
 
 // GFX90A: image_atomic_cmpswap a[6:7], v[2:5], s[12:19] dmask:0x3 unorm ; 
encoding: [0x00,0x13,0x45,0xf0,0x02,0x06,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_cmpswap a[6:7], v[2:5], s[12:19] dmask:0x3 unorm
 
 // GFX90A: image_atomic_cmpswap a[6:7], v[2:5], s[92:99] dmask:0x3 unorm ; 
encoding: [0x00,0x13,0x45,0xf0,0x02,0x06,0x17,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_cmpswap a[6:7], v[2:5], s[92:99] dmask:0x3 unorm
 
 // GFX90A: image_atomic_cmpswap a[6:9], v[2:5], s[8:15] dmask:0xf unorm ; 
encoding: [0x00,0x1f,0x45,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_cmpswap a[6:9], v[2:5], s[8:15] dmask:0xf unorm
 
 // GFX90A: image_atomic_cmpswap a[6:7], v[2:5], s[8:15] dmask:0x3 unorm glc ; 
encoding: [0x00,0x33,0x45,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_cmpswap a[6:7], v[2:5], s[8:15] dmask:0x3 unorm glc
 
 // GFX90A: image_atomic_cmpswap a[6:7], v[2:5], s[8:15] dmask:0x3 unorm slc ; 
encoding: [0x00,0x13,0x45,0xf2,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_cmpswap a[6:7], v[2:5], s[8:15] dmask:0x3 unorm slc
 
 // GFX90A: image_atomic_cmpswap a[6:7], v[2:5], s[8:15] dmask:0x3 unorm lwe ; 
encoding: [0x00,0x13,0x47,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_cmpswap a[6:7], v[2:5], s[8:15] dmask:0x3 unorm lwe
 
 // GFX90A: image_atomic_cmpswap a[6:7], v[2:5], s[8:15] dmask:0x3 unorm da ; 
encoding: [0x00,0x53,0x45,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_cmpswap a[6:7], v[2:5], s[8:15] dmask:0x3 unorm da
 
 // GFX90A: image_atomic_add a5, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x49,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_add a5, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_add a252, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x49,0xf0,0x02,0xfc,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_add a252, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_add a5, v[252:255], s[8:15] dmask:0x1 unorm ; 
encoding: [0x00,0x11,0x49,0xf0,0xfc,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_add a5, v[252:255], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_add a5, v[2:5], s[12:19] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x49,0xf0,0x02,0x05,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_add a5, v[2:5], s[12:19] dmask:0x1 unorm
 
 // GFX90A: image_atomic_add a5, v[2:5], s[92:99] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x49,0xf0,0x02,0x05,0x17,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_add a5, v[2:5], s[92:99] dmask:0x1 unorm
 
 // GFX90A: image_atomic_add a[6:7], v[2:5], s[8:15] dmask:0x3 unorm ; 
encoding: [0x00,0x13,0x49,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_add a[6:7], v[2:5], s[8:15] dmask:0x3 unorm
 
 // GFX90A: image_atomic_add a5, v[2:5], s[8:15] dmask:0x1 unorm glc ; 
encoding: [0x00,0x31,0x49,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_add a5, v[2:5], s[8:15] dmask:0x1 unorm glc
 
 // GFX90A: image_atomic_add a5, v[2:5], s[8:15] dmask:0x1 unorm slc ; 
encoding: [0x00,0x11,0x49,0xf2,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_add a5, v[2:5], s[8:15] dmask:0x1 unorm slc
 
 // GFX90A: image_atomic_add a5, v[2:5], s[8:15] dmask:0x1 unorm lwe ; 
encoding: [0x00,0x11,0x4b,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_add a5, v[2:5], s[8:15] dmask:0x1 unorm lwe
 
 // GFX90A: image_atomic_add a5, v[2:5], s[8:15] dmask:0x1 unorm da ; encoding: 
[0x00,0x51,0x49,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_add a5, v[2:5], s[8:15] dmask:0x1 unorm da
 
 // GFX90A: image_atomic_sub a5, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x4d,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_sub a5, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_sub a252, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x4d,0xf0,0x02,0xfc,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_sub a252, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_sub a5, v[252:255], s[8:15] dmask:0x1 unorm ; 
encoding: [0x00,0x11,0x4d,0xf0,0xfc,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_sub a5, v[252:255], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_sub a5, v[2:5], s[12:19] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x4d,0xf0,0x02,0x05,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_sub a5, v[2:5], s[12:19] dmask:0x1 unorm
 
 // GFX90A: image_atomic_sub a5, v[2:5], s[92:99] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x4d,0xf0,0x02,0x05,0x17,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_sub a5, v[2:5], s[92:99] dmask:0x1 unorm
 
 // GFX90A: image_atomic_sub a[6:7], v[2:5], s[8:15] dmask:0x3 unorm ; 
encoding: [0x00,0x13,0x4d,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_sub a[6:7], v[2:5], s[8:15] dmask:0x3 unorm
 
 // GFX90A: image_atomic_sub a5, v[2:5], s[8:15] dmask:0x1 unorm glc ; 
encoding: [0x00,0x31,0x4d,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_sub a5, v[2:5], s[8:15] dmask:0x1 unorm glc
 
 // GFX90A: image_atomic_sub a5, v[2:5], s[8:15] dmask:0x1 unorm slc ; 
encoding: [0x00,0x11,0x4d,0xf2,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_sub a5, v[2:5], s[8:15] dmask:0x1 unorm slc
 
 // GFX90A: image_atomic_sub a5, v[2:5], s[8:15] dmask:0x1 unorm lwe ; 
encoding: [0x00,0x11,0x4f,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_sub a5, v[2:5], s[8:15] dmask:0x1 unorm lwe
 
 // GFX90A: image_atomic_sub a5, v[2:5], s[8:15] dmask:0x1 unorm da ; encoding: 
[0x00,0x51,0x4d,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_sub a5, v[2:5], s[8:15] dmask:0x1 unorm da
 
 // GFX90A: image_atomic_smin a5, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x51,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smin a5, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_smin a252, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x51,0xf0,0x02,0xfc,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smin a252, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_smin a5, v[252:255], s[8:15] dmask:0x1 unorm ; 
encoding: [0x00,0x11,0x51,0xf0,0xfc,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smin a5, v[252:255], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_smin a5, v[2:5], s[12:19] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x51,0xf0,0x02,0x05,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smin a5, v[2:5], s[12:19] dmask:0x1 unorm
 
 // GFX90A: image_atomic_smin a5, v[2:5], s[92:99] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x51,0xf0,0x02,0x05,0x17,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smin a5, v[2:5], s[92:99] dmask:0x1 unorm
 
 // GFX90A: image_atomic_smin a[6:7], v[2:5], s[8:15] dmask:0x3 unorm ; 
encoding: [0x00,0x13,0x51,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smin a[6:7], v[2:5], s[8:15] dmask:0x3 unorm
 
 // GFX90A: image_atomic_smin a5, v[2:5], s[8:15] dmask:0x1 unorm glc ; 
encoding: [0x00,0x31,0x51,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smin a5, v[2:5], s[8:15] dmask:0x1 unorm glc
 
 // GFX90A: image_atomic_smin a5, v[2:5], s[8:15] dmask:0x1 unorm slc ; 
encoding: [0x00,0x11,0x51,0xf2,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smin a5, v[2:5], s[8:15] dmask:0x1 unorm slc
 
 // GFX90A: image_atomic_smin a5, v[2:5], s[8:15] dmask:0x1 unorm lwe ; 
encoding: [0x00,0x11,0x53,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smin a5, v[2:5], s[8:15] dmask:0x1 unorm lwe
 
 // GFX90A: image_atomic_smin a5, v[2:5], s[8:15] dmask:0x1 unorm da ; 
encoding: [0x00,0x51,0x51,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smin a5, v[2:5], s[8:15] dmask:0x1 unorm da
 
 // GFX90A: image_atomic_umin a5, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x55,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umin a5, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_umin a252, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x55,0xf0,0x02,0xfc,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umin a252, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_umin a5, v[252:255], s[8:15] dmask:0x1 unorm ; 
encoding: [0x00,0x11,0x55,0xf0,0xfc,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umin a5, v[252:255], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_umin a5, v[2:5], s[12:19] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x55,0xf0,0x02,0x05,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umin a5, v[2:5], s[12:19] dmask:0x1 unorm
 
 // GFX90A: image_atomic_umin a5, v[2:5], s[92:99] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x55,0xf0,0x02,0x05,0x17,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umin a5, v[2:5], s[92:99] dmask:0x1 unorm
 
 // GFX90A: image_atomic_umin a[6:7], v[2:5], s[8:15] dmask:0x3 unorm ; 
encoding: [0x00,0x13,0x55,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umin a[6:7], v[2:5], s[8:15] dmask:0x3 unorm
 
 // GFX90A: image_atomic_umin a5, v[2:5], s[8:15] dmask:0x1 unorm glc ; 
encoding: [0x00,0x31,0x55,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umin a5, v[2:5], s[8:15] dmask:0x1 unorm glc
 
 // GFX90A: image_atomic_umin a5, v[2:5], s[8:15] dmask:0x1 unorm slc ; 
encoding: [0x00,0x11,0x55,0xf2,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umin a5, v[2:5], s[8:15] dmask:0x1 unorm slc
 
 // GFX90A: image_atomic_umin a5, v[2:5], s[8:15] dmask:0x1 unorm lwe ; 
encoding: [0x00,0x11,0x57,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umin a5, v[2:5], s[8:15] dmask:0x1 unorm lwe
 
 // GFX90A: image_atomic_umin a5, v[2:5], s[8:15] dmask:0x1 unorm da ; 
encoding: [0x00,0x51,0x55,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umin a5, v[2:5], s[8:15] dmask:0x1 unorm da
 
 // GFX90A: image_atomic_smax a5, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x59,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smax a5, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_smax a252, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x59,0xf0,0x02,0xfc,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smax a252, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_smax a5, v[252:255], s[8:15] dmask:0x1 unorm ; 
encoding: [0x00,0x11,0x59,0xf0,0xfc,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smax a5, v[252:255], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_smax a5, v[2:5], s[12:19] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x59,0xf0,0x02,0x05,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smax a5, v[2:5], s[12:19] dmask:0x1 unorm
 
 // GFX90A: image_atomic_smax a5, v[2:5], s[92:99] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x59,0xf0,0x02,0x05,0x17,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smax a5, v[2:5], s[92:99] dmask:0x1 unorm
 
 // GFX90A: image_atomic_smax a[6:7], v[2:5], s[8:15] dmask:0x3 unorm ; 
encoding: [0x00,0x13,0x59,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smax a[6:7], v[2:5], s[8:15] dmask:0x3 unorm
 
 // GFX90A: image_atomic_smax a5, v[2:5], s[8:15] dmask:0x1 unorm glc ; 
encoding: [0x00,0x31,0x59,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smax a5, v[2:5], s[8:15] dmask:0x1 unorm glc
 
 // GFX90A: image_atomic_smax a5, v[2:5], s[8:15] dmask:0x1 unorm slc ; 
encoding: [0x00,0x11,0x59,0xf2,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smax a5, v[2:5], s[8:15] dmask:0x1 unorm slc
 
 // GFX90A: image_atomic_smax a5, v[2:5], s[8:15] dmask:0x1 unorm lwe ; 
encoding: [0x00,0x11,0x5b,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smax a5, v[2:5], s[8:15] dmask:0x1 unorm lwe
 
 // GFX90A: image_atomic_smax a5, v[2:5], s[8:15] dmask:0x1 unorm da ; 
encoding: [0x00,0x51,0x59,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_smax a5, v[2:5], s[8:15] dmask:0x1 unorm da
 
 // GFX90A: image_atomic_umax a5, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x5d,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umax a5, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_umax a252, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x5d,0xf0,0x02,0xfc,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umax a252, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_umax a5, v[252:255], s[8:15] dmask:0x1 unorm ; 
encoding: [0x00,0x11,0x5d,0xf0,0xfc,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umax a5, v[252:255], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_umax a5, v[2:5], s[12:19] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x5d,0xf0,0x02,0x05,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umax a5, v[2:5], s[12:19] dmask:0x1 unorm
 
 // GFX90A: image_atomic_umax a5, v[2:5], s[92:99] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x5d,0xf0,0x02,0x05,0x17,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umax a5, v[2:5], s[92:99] dmask:0x1 unorm
 
 // GFX90A: image_atomic_umax a[6:7], v[2:5], s[8:15] dmask:0x3 unorm ; 
encoding: [0x00,0x13,0x5d,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umax a[6:7], v[2:5], s[8:15] dmask:0x3 unorm
 
 // GFX90A: image_atomic_umax a5, v[2:5], s[8:15] dmask:0x1 unorm glc ; 
encoding: [0x00,0x31,0x5d,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umax a5, v[2:5], s[8:15] dmask:0x1 unorm glc
 
 // GFX90A: image_atomic_umax a5, v[2:5], s[8:15] dmask:0x1 unorm slc ; 
encoding: [0x00,0x11,0x5d,0xf2,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umax a5, v[2:5], s[8:15] dmask:0x1 unorm slc
 
 // GFX90A: image_atomic_umax a5, v[2:5], s[8:15] dmask:0x1 unorm lwe ; 
encoding: [0x00,0x11,0x5f,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umax a5, v[2:5], s[8:15] dmask:0x1 unorm lwe
 
 // GFX90A: image_atomic_umax a5, v[2:5], s[8:15] dmask:0x1 unorm da ; 
encoding: [0x00,0x51,0x5d,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_umax a5, v[2:5], s[8:15] dmask:0x1 unorm da
 
 // GFX90A: image_atomic_and a5, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x61,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_and a5, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_and a252, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x61,0xf0,0x02,0xfc,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_and a252, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_and a5, v[252:255], s[8:15] dmask:0x1 unorm ; 
encoding: [0x00,0x11,0x61,0xf0,0xfc,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_and a5, v[252:255], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_and a5, v[2:5], s[12:19] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x61,0xf0,0x02,0x05,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_and a5, v[2:5], s[12:19] dmask:0x1 unorm
 
 // GFX90A: image_atomic_and a5, v[2:5], s[92:99] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x61,0xf0,0x02,0x05,0x17,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_and a5, v[2:5], s[92:99] dmask:0x1 unorm
 
 // GFX90A: image_atomic_and a[6:7], v[2:5], s[8:15] dmask:0x3 unorm ; 
encoding: [0x00,0x13,0x61,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_and a[6:7], v[2:5], s[8:15] dmask:0x3 unorm
 
 // GFX90A: image_atomic_and a5, v[2:5], s[8:15] dmask:0x1 unorm glc ; 
encoding: [0x00,0x31,0x61,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_and a5, v[2:5], s[8:15] dmask:0x1 unorm glc
 
 // GFX90A: image_atomic_and a5, v[2:5], s[8:15] dmask:0x1 unorm slc ; 
encoding: [0x00,0x11,0x61,0xf2,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_and a5, v[2:5], s[8:15] dmask:0x1 unorm slc
 
 // GFX90A: image_atomic_and a5, v[2:5], s[8:15] dmask:0x1 unorm lwe ; 
encoding: [0x00,0x11,0x63,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_and a5, v[2:5], s[8:15] dmask:0x1 unorm lwe
 
 // GFX90A: image_atomic_and a5, v[2:5], s[8:15] dmask:0x1 unorm da ; encoding: 
[0x00,0x51,0x61,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_and a5, v[2:5], s[8:15] dmask:0x1 unorm da
 
 // GFX90A: image_atomic_or a5, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x65,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_or a5, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_or a252, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x65,0xf0,0x02,0xfc,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_or a252, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_or a5, v[252:255], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x65,0xf0,0xfc,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_or a5, v[252:255], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_or a5, v[2:5], s[12:19] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x65,0xf0,0x02,0x05,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_or a5, v[2:5], s[12:19] dmask:0x1 unorm
 
 // GFX90A: image_atomic_or a5, v[2:5], s[92:99] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x65,0xf0,0x02,0x05,0x17,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_or a5, v[2:5], s[92:99] dmask:0x1 unorm
 
 // GFX90A: image_atomic_or a[6:7], v[2:5], s[8:15] dmask:0x3 unorm ; encoding: 
[0x00,0x13,0x65,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_or a[6:7], v[2:5], s[8:15] dmask:0x3 unorm
 
 // GFX90A: image_atomic_or a5, v[2:5], s[8:15] dmask:0x1 unorm glc ; encoding: 
[0x00,0x31,0x65,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_or a5, v[2:5], s[8:15] dmask:0x1 unorm glc
 
 // GFX90A: image_atomic_or a5, v[2:5], s[8:15] dmask:0x1 unorm slc ; encoding: 
[0x00,0x11,0x65,0xf2,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_or a5, v[2:5], s[8:15] dmask:0x1 unorm slc
 
 // GFX90A: image_atomic_or a5, v[2:5], s[8:15] dmask:0x1 unorm lwe ; encoding: 
[0x00,0x11,0x67,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_or a5, v[2:5], s[8:15] dmask:0x1 unorm lwe
 
 // GFX90A: image_atomic_or a5, v[2:5], s[8:15] dmask:0x1 unorm da ; encoding: 
[0x00,0x51,0x65,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_or a5, v[2:5], s[8:15] dmask:0x1 unorm da
 
 // GFX90A: image_atomic_xor a5, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x69,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_xor a5, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_xor a252, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x69,0xf0,0x02,0xfc,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_xor a252, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_xor a5, v[252:255], s[8:15] dmask:0x1 unorm ; 
encoding: [0x00,0x11,0x69,0xf0,0xfc,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_xor a5, v[252:255], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_xor a5, v[2:5], s[12:19] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x69,0xf0,0x02,0x05,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_xor a5, v[2:5], s[12:19] dmask:0x1 unorm
 
 // GFX90A: image_atomic_xor a5, v[2:5], s[92:99] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x69,0xf0,0x02,0x05,0x17,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_xor a5, v[2:5], s[92:99] dmask:0x1 unorm
 
 // GFX90A: image_atomic_xor a[6:7], v[2:5], s[8:15] dmask:0x3 unorm ; 
encoding: [0x00,0x13,0x69,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_xor a[6:7], v[2:5], s[8:15] dmask:0x3 unorm
 
 // GFX90A: image_atomic_xor a5, v[2:5], s[8:15] dmask:0x1 unorm glc ; 
encoding: [0x00,0x31,0x69,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_xor a5, v[2:5], s[8:15] dmask:0x1 unorm glc
 
 // GFX90A: image_atomic_xor a5, v[2:5], s[8:15] dmask:0x1 unorm slc ; 
encoding: [0x00,0x11,0x69,0xf2,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_xor a5, v[2:5], s[8:15] dmask:0x1 unorm slc
 
 // GFX90A: image_atomic_xor a5, v[2:5], s[8:15] dmask:0x1 unorm lwe ; 
encoding: [0x00,0x11,0x6b,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_xor a5, v[2:5], s[8:15] dmask:0x1 unorm lwe
 
 // GFX90A: image_atomic_xor a5, v[2:5], s[8:15] dmask:0x1 unorm da ; encoding: 
[0x00,0x51,0x69,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_xor a5, v[2:5], s[8:15] dmask:0x1 unorm da
 
 // GFX90A: image_atomic_inc a5, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x6d,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_inc a5, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_inc a252, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x6d,0xf0,0x02,0xfc,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_inc a252, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_inc a5, v[252:255], s[8:15] dmask:0x1 unorm ; 
encoding: [0x00,0x11,0x6d,0xf0,0xfc,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_inc a5, v[252:255], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_inc a5, v[2:5], s[12:19] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x6d,0xf0,0x02,0x05,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_inc a5, v[2:5], s[12:19] dmask:0x1 unorm
 
 // GFX90A: image_atomic_inc a5, v[2:5], s[92:99] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x6d,0xf0,0x02,0x05,0x17,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_inc a5, v[2:5], s[92:99] dmask:0x1 unorm
 
 // GFX90A: image_atomic_inc a[6:7], v[2:5], s[8:15] dmask:0x3 unorm ; 
encoding: [0x00,0x13,0x6d,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_inc a[6:7], v[2:5], s[8:15] dmask:0x3 unorm
 
 // GFX90A: image_atomic_inc a5, v[2:5], s[8:15] dmask:0x1 unorm glc ; 
encoding: [0x00,0x31,0x6d,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_inc a5, v[2:5], s[8:15] dmask:0x1 unorm glc
 
 // GFX90A: image_atomic_inc a5, v[2:5], s[8:15] dmask:0x1 unorm slc ; 
encoding: [0x00,0x11,0x6d,0xf2,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_inc a5, v[2:5], s[8:15] dmask:0x1 unorm slc
 
 // GFX90A: image_atomic_inc a5, v[2:5], s[8:15] dmask:0x1 unorm lwe ; 
encoding: [0x00,0x11,0x6f,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_inc a5, v[2:5], s[8:15] dmask:0x1 unorm lwe
 
 // GFX90A: image_atomic_inc a5, v[2:5], s[8:15] dmask:0x1 unorm da ; encoding: 
[0x00,0x51,0x6d,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_inc a5, v[2:5], s[8:15] dmask:0x1 unorm da
 
 // GFX90A: image_atomic_dec a5, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x71,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_dec a5, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_dec a252, v[2:5], s[8:15] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x71,0xf0,0x02,0xfc,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_dec a252, v[2:5], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_dec a5, v[252:255], s[8:15] dmask:0x1 unorm ; 
encoding: [0x00,0x11,0x71,0xf0,0xfc,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_dec a5, v[252:255], s[8:15] dmask:0x1 unorm
 
 // GFX90A: image_atomic_dec a5, v[2:5], s[12:19] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x71,0xf0,0x02,0x05,0x03,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_dec a5, v[2:5], s[12:19] dmask:0x1 unorm
 
 // GFX90A: image_atomic_dec a5, v[2:5], s[92:99] dmask:0x1 unorm ; encoding: 
[0x00,0x11,0x71,0xf0,0x02,0x05,0x17,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_dec a5, v[2:5], s[92:99] dmask:0x1 unorm
 
 // GFX90A: image_atomic_dec a[6:7], v[2:5], s[8:15] dmask:0x3 unorm ; 
encoding: [0x00,0x13,0x71,0xf0,0x02,0x06,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_dec a[6:7], v[2:5], s[8:15] dmask:0x3 unorm
 
 // GFX90A: image_atomic_dec a5, v[2:5], s[8:15] dmask:0x1 unorm glc ; 
encoding: [0x00,0x31,0x71,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_dec a5, v[2:5], s[8:15] dmask:0x1 unorm glc
 
 // GFX90A: image_atomic_dec a5, v[2:5], s[8:15] dmask:0x1 unorm slc ; 
encoding: [0x00,0x11,0x71,0xf2,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_dec a5, v[2:5], s[8:15] dmask:0x1 unorm slc
 
 // GFX90A: image_atomic_dec a5, v[2:5], s[8:15] dmask:0x1 unorm lwe ; 
encoding: [0x00,0x11,0x73,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_dec a5, v[2:5], s[8:15] dmask:0x1 unorm lwe
 
 // GFX90A: image_atomic_dec a5, v[2:5], s[8:15] dmask:0x1 unorm da ; encoding: 
[0x00,0x51,0x71,0xf0,0x02,0x05,0x02,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_atomic_dec a5, v[2:5], s[8:15] dmask:0x1 unorm da
 
 // GFX90A: image_sample a5, v[0:3], s[8:15], s[12:15] dmask:0x1 ; encoding: 
[0x00,0x01,0x81,0xf0,0x00,0x05,0x62,0x00]
-// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: operands are not valid for this 
GPU or mode
+// NOT-GFX90A: :[[@LINE+1]]:{{[0-9]+}}: error: invalid register class: agpr 
loads and stores not supported on this GPU
 image_sample a5, v[0:3], s[8:15], s[12:15] dmask:0x1
diff --git a/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s 
b/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
index 0409a98f0f77e..c935c378bbba2 100644
--- a/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
+++ b/llvm/test/MC/AMDGPU/misaligned-vgpr-tuples-err.s
@@ -32,13 +32,13 @@ image_load v[1:4], v2, s[0:7] dmask:0xf unorm
 // GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU 
or mode
 
 image_load a[1:2], v2, s[0:7] dmask:0x3 unorm
-// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction
+// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU 
or mode
 
 image_load a[1:3], v2, s[0:7] dmask:0x7 unorm
-// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction
+// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU 
or mode
 
 image_load a[1:4], v2, s[0:7] dmask:0xf unorm
-// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction
+// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU 
or mode
 
 
 image_store v[193:194], v[238:241], s[28:35] dmask:0x3 unorm
@@ -51,13 +51,13 @@ image_store v[193:196], v[238:241], s[28:35] dmask:0xf unorm
 // GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU 
or mode
 
 image_store a[193:194], v[238:241], s[28:35] dmask:0x3 unorm
-// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction
+// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU 
or mode
 
 image_store a[193:195], v[238:241], s[28:35] dmask:0x7 unorm
-// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction
+// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU 
or mode
 
 image_store a[193:196], v[238:241], s[28:35] dmask:0xf unorm
-// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction
+// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU 
or mode
 
 
 image_atomic_swap v4, v[193:196], s[28:35] dmask:0x1 unorm glc
@@ -81,16 +81,16 @@ image_atomic_cmpswap v[4:7], v[193:196], s[28:35] dmask:0xf 
unorm glc
 
 
 image_atomic_cmpswap a[5:6], v[192:195], s[28:35] dmask:0x3 unorm glc
-// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction
+// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU 
or mode
 
 image_atomic_cmpswap a[4:5], v[193:196], s[28:35] dmask:0x3 unorm glc
-// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction
+// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU 
or mode
 
 image_atomic_cmpswap a[5:8], v[192:195], s[28:35] dmask:0xf unorm glc
-// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction
+// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU 
or mode
 
 image_atomic_cmpswap a[4:7], v[193:196], s[28:35] dmask:0xf unorm glc
-// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction
+// GFX90A: :[[@LINE-1]]:{{[0-9]+}}: error: operands are not valid for this GPU 
or mode
 
 
 v_mfma_f32_32x32x8f16 a[0:15], a[1:2], v[0:1], a[0:15]

_______________________________________________
llvm-branch-commits mailing list
llvm-branch-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits