llvm/test/MC/AMDGPU/gfx8_unsupported.s

   1 // RUN: not llvm-mc -triple=amdgcn -mcpu=tonga %s 2>&1 | FileCheck --implicit-check-not=error: %s
   2
   3 //===----------------------------------------------------------------------===//
   4 // Unsupported instructions.
   5 //===----------------------------------------------------------------------===//
   6
   7 image_sample_c_cd_cl_g16 v[5:6], v[1:5], s[8:15], s[12:15] dmask:0x3
   8 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
   9
  10 image_sample_c_cd_cl_o_g16 v[5:6], v[1:6], s[8:15], s[12:15] dmask:0x3
  11 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  12
  13 image_sample_c_cd_g16 v[5:6], v[1:4], s[8:15], s[12:15] dmask:0x3
  14 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  15
  16 image_sample_c_cd_o_g16 v[5:6], v[1:5], s[8:15], s[12:15] dmask:0x3
  17 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  18
  19 image_sample_c_d_cl_g16 v[5:6], v[1:5], s[8:15], s[12:15] dmask:0x3
  20 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  21
  22 image_sample_c_d_cl_o_g16 v[5:6], v[1:6], s[8:15], s[12:15] dmask:0x3
  23 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  24
  25 image_sample_c_d_g16 v[5:6], v[1:4], s[8:15], s[12:15] dmask:0x3
  26 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  27
  28 image_sample_c_d_o_g16 v[5:6], v[1:5], s[8:15], s[12:15] dmask:0x3
  29 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  30
  31 image_sample_cd_cl_g16 v[5:6], v[1:4], s[8:15], s[12:15] dmask:0x3
  32 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  33
  34 image_sample_cd_cl_o_g16 v[5:6], v[1:5], s[8:15], s[12:15] dmask:0x3
  35 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  36
  37 image_sample_cd_g16 v[5:6], v[1:3], s[8:15], s[12:15] dmask:0x3
  38 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  39
  40 image_sample_cd_o_g16 v[5:6], v[1:4], s[8:15], s[12:15] dmask:0x3
  41 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  42
  43 image_sample_d_cl_g16 v[5:6], v[1:4], s[8:15], s[12:15] dmask:0x3
  44 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  45
  46 image_sample_d_cl_o_g16 v[5:6], v[1:5], s[8:15], s[12:15] dmask:0x3
  47 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  48
  49 image_sample_d_g16 v[5:6], v[1:3], s[8:15], s[12:15] dmask:0x3
  50 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  51
  52 image_sample_d_o_g16 v[5:6], v[1:4], s[8:15], s[12:15] dmask:0x3
  53 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  54
  55 image_gather4h v[251:254], v[1:2], s[8:15], s[12:15] dmask:0x1
  56 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  57
  58 buffer_atomic_add_f32 v255, off, s[8:11], s3 offset:4095
  59 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  60
  61 buffer_atomic_fcmpswap v[0:1], off, s[0:3], s0 offset:4095
  62 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  63
  64 buffer_atomic_fcmpswap_x2 v[0:3], off, s[0:3], s0 offset:4095
  65 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  66
  67 buffer_atomic_fmax v0, off, s[0:3], s0 offset:4095 glc
  68 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  69
  70 buffer_atomic_fmax_x2 v[0:1], v0, s[0:3], s0 idxen offset:4095
  71 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  72
  73 buffer_atomic_fmin v0, off, s[0:3], s0
  74 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  75
  76 buffer_atomic_fmin_x2 v[0:1], off, s[0:3], s0 offset:4095 slc
  77 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  78
  79 buffer_atomic_pk_add_f16 v255, off, s[8:11], s3 offset:4095
  80 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  81
  82 buffer_gl0_inv
  83 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  84
  85 buffer_gl1_inv
  86 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  87
  88 buffer_load_format_d16_hi_x v5, off, s[8:11], s3
  89 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  90
  91 buffer_load_sbyte_d16 v1, off, s[4:7], s1
  92 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  93
  94 buffer_load_sbyte_d16_hi v1, off, s[4:7], s1
  95 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  96
  97 buffer_load_short_d16 v1, off, s[4:7], s1
  98 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  99
 100 buffer_load_short_d16_hi v1, off, s[4:7], s1
 101 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 102
 103 buffer_load_ubyte_d16 v1, off, s[4:7], s1
 104 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 105
 106 buffer_load_ubyte_d16_hi v1, off, s[4:7], s1
 107 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 108
 109 buffer_store_byte_d16_hi v1, off, s[12:15], -1 offset:4095
 110 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 111
 112 buffer_store_format_d16_hi_x v1, off, s[12:15], s4 offset:4095 glc
 113 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 114
 115 buffer_store_short_d16_hi v1, off, s[12:15], -1 offset:4095
 116 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 117
 118 ds_read_addtid_b32 v255 offset:65535
 119 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 120
 121 ds_read_i8_d16 v255, v1 offset:65535
 122 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 123
 124 ds_read_i8_d16_hi v255, v1 offset:65535
 125 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 126
 127 ds_read_u16_d16 v255, v1 offset:65535
 128 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 129
 130 ds_read_u16_d16_hi v255, v1 offset:65535
 131 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 132
 133 ds_read_u8_d16 v255, v1 offset:65535
 134 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 135
 136 ds_read_u8_d16_hi v255, v1 offset:65535
 137 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 138
 139 ds_write_addtid_b32 v255 offset:65535
 140 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 141
 142 ds_write_b16_d16_hi v1, v2
 143 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 144
 145 ds_write_b8_d16_hi v1, v2
 146 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 147
 148 flat_atomic_fcmpswap v0, v[1:2], v[2:3] glc
 149 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 150
 151 flat_atomic_fcmpswap_x2 v[0:1], v[1:2], v[2:5] glc
 152 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 153
 154 flat_atomic_fmax v0, v[1:2], v2 glc
 155 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 156
 157 flat_atomic_fmax_x2 v[0:1], v[1:2], v[2:3] glc
 158 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 159
 160 flat_atomic_fmin v0, v[1:2], v2 glc
 161 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 162
 163 flat_atomic_fmin_x2 v[0:1], v[1:2], v[2:3] glc
 164 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 165
 166 flat_load_sbyte_d16 v1, v[3:4]
 167 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 168
 169 flat_load_sbyte_d16_hi v1, v[3:4]
 170 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 171
 172 flat_load_short_d16 v1, v[3:4]
 173 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 174
 175 flat_load_short_d16_hi v1, v[3:4]
 176 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 177
 178 flat_load_ubyte_d16 v1, v[3:4]
 179 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 180
 181 flat_load_ubyte_d16_hi v1, v[3:4]
 182 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 183
 184 flat_store_byte_d16_hi v[1:2], v2
 185 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 186
 187 flat_store_short_d16_hi v[1:2], v2
 188 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 189
 190 global_atomic_add v0, v[1:2], v2, off glc slc
 191 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 192
 193 global_atomic_add_f32 v[1:2], v2, off
 194 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 195
 196 global_atomic_add_x2 v[1:2], v[254:255], off offset:-1
 197 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 198
 199 global_atomic_and v[1:2], v2, off
 200 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 201
 202 global_atomic_and_x2 v[1:2], v[254:255], off offset:-1
 203 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 204
 205 global_atomic_cmpswap v[1:2], v[254:255], off offset:-1
 206 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 207
 208 global_atomic_cmpswap_x2 v[1:2], v[252:255], off offset:-1
 209 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 210
 211 global_atomic_dec v[1:2], v2, off
 212 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 213
 214 global_atomic_dec_x2 v[1:2], v[254:255], off offset:-1
 215 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 216
 217 global_atomic_inc v[1:2], v2, off
 218 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 219
 220 global_atomic_inc_x2 v[1:2], v[254:255], off offset:-1
 221 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 222
 223 global_atomic_or v[1:2], v2, off
 224 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 225
 226 global_atomic_or_x2 v[1:2], v[254:255], off offset:-1
 227 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 228
 229 global_atomic_pk_add_f16 v[1:2], v2, off
 230 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 231
 232 global_atomic_smax v[1:2], v2, off
 233 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 234
 235 global_atomic_smax_x2 v[1:2], v[254:255], off offset:-1
 236 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 237
 238 global_atomic_smin v[1:2], v2, off
 239 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 240
 241 global_atomic_smin_x2 v[1:2], v[254:255], off offset:-1
 242 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 243
 244 global_atomic_sub v[1:2], v2, off
 245 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 246
 247 global_atomic_sub_x2 v[1:2], v[254:255], off offset:-1
 248 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 249
 250 global_atomic_swap v[1:2], v2, off
 251 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 252
 253 global_atomic_swap_x2 v[1:2], v[254:255], off offset:-1
 254 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 255
 256 global_atomic_umax v[1:2], v2, off
 257 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 258
 259 global_atomic_umax_x2 v[1:2], v[254:255], off offset:-1
 260 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 261
 262 global_atomic_umin v[1:2], v2, off
 263 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 264
 265 global_atomic_umin_x2 v[1:2], v[254:255], off offset:-1
 266 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 267
 268 global_atomic_xor v[1:2], v2, off
 269 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 270
 271 global_atomic_xor_x2 v[1:2], v[254:255], off offset:-1
 272 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 273
 274 global_load_dword v1, v3, s[2:3]
 275 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 276
 277 global_load_dwordx2 v[1:2], v[3:4], off
 278 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 279
 280 global_load_dwordx3 v[1:3], v[3:4], off
 281 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 282
 283 global_load_dwordx4 v[1:4], v[3:4], off
 284 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 285
 286 global_load_sbyte v1, v[3:4], off
 287 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 288
 289 global_load_sbyte_d16 v1, v[3:4], off
 290 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 291
 292 global_load_sbyte_d16_hi v1, v[3:4], off
 293 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 294
 295 global_load_short_d16 v1, v[3:4], off
 296 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 297
 298 global_load_short_d16_hi v1, v[3:4], off
 299 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 300
 301 global_load_sshort v1, v[3:4], off
 302 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 303
 304 global_load_ubyte v1, v[3:4], off
 305 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 306
 307 global_load_ubyte_d16 v1, v[3:4], off
 308 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 309
 310 global_load_ubyte_d16_hi v1, v[3:4], off
 311 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 312
 313 global_load_ushort v1, v[3:4], off
 314 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 315
 316 global_store_byte v[1:2], v2, off
 317 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 318
 319 global_store_byte_d16_hi v[1:2], v2, off
 320 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 321
 322 global_store_dword v254, v1, s[2:3] offset:16
 323 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 324
 325 global_store_dwordx2 v[1:2], v[254:255], off offset:-1
 326 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 327
 328 global_store_dwordx3 v[1:2], v[253:255], off offset:-1
 329 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 330
 331 global_store_dwordx4 v[1:2], v[252:255], off offset:-1
 332 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 333
 334 global_store_short v[1:2], v2, off
 335 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 336
 337 global_store_short_d16_hi v[1:2], v2, off
 338 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 339
 340 s_and_saveexec_b32 exec_hi, s1
 341 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 342
 343 s_andn1_saveexec_b32 exec_hi, s1
 344 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 345
 346 s_andn1_saveexec_b64 exec, s[2:3]
 347 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 348
 349 s_andn1_wrexec_b32 exec_hi, s1
 350 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 351
 352 s_andn1_wrexec_b64 exec, s[2:3]
 353 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 354
 355 s_andn2_saveexec_b32 exec_hi, s1
 356 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 357
 358 s_andn2_wrexec_b32 exec_hi, s1
 359 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 360
 361 s_andn2_wrexec_b64 exec, s[2:3]
 362 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 363
 364 s_atomic_add s5, s[2:3], 0x0
 365 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 366
 367 s_atomic_add_x2 s[10:11], s[2:3], s101
 368 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 369
 370 s_atomic_and s5, s[2:3], s101
 371 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 372
 373 s_atomic_and_x2 s[10:11], s[2:3], 0x0
 374 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 375
 376 s_atomic_cmpswap s[10:11], s[2:3], 0x0
 377 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 378
 379 s_atomic_cmpswap_x2 s[20:23], s[2:3], 0x0
 380 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 381
 382 s_atomic_dec s5, s[2:3], s0 glc
 383 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 384
 385 s_atomic_dec_x2 s[10:11], s[2:3], s101
 386 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 387
 388 s_atomic_inc s5, s[2:3], s0 glc
 389 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 390
 391 s_atomic_inc_x2 s[10:11], s[2:3], s101
 392 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 393
 394 s_atomic_or s5, s[2:3], 0x0
 395 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 396
 397 s_atomic_or_x2 s[10:11], s[2:3], s0 glc
 398 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 399
 400 s_atomic_smax s5, s[2:3], s101
 401 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 402
 403 s_atomic_smax_x2 s[10:11], s[2:3], s0 glc
 404 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 405
 406 s_atomic_smin s5, s[2:3], s101
 407 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 408
 409 s_atomic_smin_x2 s[10:11], s[2:3], s0 glc
 410 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 411
 412 s_atomic_sub s5, s[2:3], s101
 413 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 414
 415 s_atomic_sub_x2 s[10:11], s[2:3], s0 glc
 416 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 417
 418 s_atomic_swap s5, s[2:3], -1
 419 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 420
 421 s_atomic_swap_x2 s[10:11], s[2:3], s0 glc
 422 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 423
 424 s_atomic_umax s5, s[2:3], s0 glc
 425 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 426
 427 s_atomic_umax_x2 s[10:11], s[2:3], s101
 428 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 429
 430 s_atomic_umin s5, s[2:3], s101
 431 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 432
 433 s_atomic_umin_x2 s[10:11], s[2:3], s0 glc
 434 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 435
 436 s_atomic_xor s5, s[2:3], s101
 437 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 438
 439 s_atomic_xor_x2 s[10:11], s[2:3], s0 glc
 440 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 441
 442 s_bitreplicate_b64_b32 exec, s2
 443 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 444
 445 s_buffer_atomic_add s5, s[4:7], 0x0
 446 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 447
 448 s_buffer_atomic_add_x2 s[10:11], s[4:7], s0
 449 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 450
 451 s_buffer_atomic_and s101, s[4:7], s0
 452 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 453
 454 s_buffer_atomic_and_x2 s[10:11], s[8:11], s0
 455 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 456
 457 s_buffer_atomic_cmpswap s[10:11], s[4:7], 0x0
 458 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 459
 460 s_buffer_atomic_cmpswap_x2 s[20:23], s[4:7], 0x0
 461 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 462
 463 s_buffer_atomic_dec s5, s[4:7], s0
 464 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 465
 466 s_buffer_atomic_dec_x2 s[10:11], s[4:7], s0 glc
 467 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 468
 469 s_buffer_atomic_inc s101, s[4:7], s0
 470 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 471
 472 s_buffer_atomic_inc_x2 s[10:11], s[4:7], 0x0
 473 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 474
 475 s_buffer_atomic_or s5, s[8:11], s0
 476 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 477
 478 s_buffer_atomic_or_x2 s[10:11], s[96:99], s0
 479 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 480
 481 s_buffer_atomic_smax s5, s[4:7], s101
 482 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 483
 484 s_buffer_atomic_smax_x2 s[100:101], s[4:7], s0
 485 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 486
 487 s_buffer_atomic_smin s5, s[4:7], 0x0
 488 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 489
 490 s_buffer_atomic_smin_x2 s[12:13], s[4:7], s0
 491 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 492
 493 s_buffer_atomic_sub s5, s[4:7], s0 glc
 494 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 495
 496 s_buffer_atomic_sub_x2 s[10:11], s[4:7], s0
 497 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 498
 499 s_buffer_atomic_swap s5, s[4:7], -1
 500 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 501
 502 s_buffer_atomic_swap_x2 s[10:11], s[4:7], s0 glc
 503 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 504
 505 s_buffer_atomic_umax s5, s[4:7], s0
 506 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 507
 508 s_buffer_atomic_umax_x2 s[10:11], s[4:7], s0 glc
 509 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 510
 511 s_buffer_atomic_umin s5, s[4:7], s0
 512 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 513
 514 s_buffer_atomic_umin_x2 s[10:11], s[4:7], s0 glc
 515 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 516
 517 s_buffer_atomic_xor s5, s[4:7], s0
 518 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 519
 520 s_buffer_atomic_xor_x2 s[10:11], s[4:7], s0 glc
 521 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 522
 523 s_call_b64 exec, 0x1234
 524 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 525
 526 s_clause 0x0
 527 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 528
 529 s_code_end
 530 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 531
 532 s_dcache_discard s[2:3], 0x0
 533 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 534
 535 s_dcache_discard_x2 s[2:3], 0x0
 536 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 537
 538 s_denorm_mode 0x0
 539 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 540
 541 s_endpgm_ordered_ps_done
 542 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 543
 544 s_get_waveid_in_workgroup s0
 545 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 546
 547 s_gl1_inv
 548 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 549
 550 s_inst_prefetch 0x0
 551 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 552
 553 s_lshl1_add_u32 exec_hi, s1, s2
 554 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 555
 556 s_lshl2_add_u32 exec_hi, s1, s2
 557 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 558
 559 s_lshl3_add_u32 exec_hi, s1, s2
 560 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 561
 562 s_lshl4_add_u32 exec_hi, s1, s2
 563 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 564
 565 s_movrelsd_2_b32 s0, s1
 566 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 567
 568 s_mul_hi_i32 exec_hi, s1, s2
 569 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 570
 571 s_mul_hi_u32 exec_hi, s1, s2
 572 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 573
 574 s_nand_saveexec_b32 exec_hi, s1
 575 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 576
 577 s_nor_saveexec_b32 exec_hi, s1
 578 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 579
 580 s_or_saveexec_b32 exec_hi, s1
 581 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 582
 583 s_orn1_saveexec_b32 exec_hi, s1
 584 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 585
 586 s_orn1_saveexec_b64 exec, s[2:3]
 587 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 588
 589 s_orn2_saveexec_b32 exec_hi, s1
 590 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 591
 592 s_pack_hh_b32_b16 exec_hi, s1, s2
 593 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 594
 595 s_pack_lh_b32_b16 exec_hi, s1, s2
 596 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 597
 598 s_pack_ll_b32_b16 exec_hi, s1, s2
 599 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 600
 601 s_round_mode 0x0
 602 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 603
 604 s_scratch_load_dword s5, s[2:3], s0 glc
 605 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 606
 607 s_scratch_load_dwordx2 s[100:101], s[2:3], s0
 608 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 609
 610 s_scratch_load_dwordx4 s[20:23], s[4:5], s0
 611 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 612
 613 s_scratch_store_dword s1, s[4:5], 0x123 glc
 614 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 615
 616 s_scratch_store_dwordx2 s[2:3], s[4:5], s101 glc
 617 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 618
 619 s_scratch_store_dwordx4 s[4:7], s[4:5], s0 glc
 620 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 621
 622 s_subvector_loop_begin exec_hi, 0x1234
 623 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 624
 625 s_subvector_loop_end exec_hi, 0x1234
 626 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 627
 628 s_ttracedata_imm 0x0
 629 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 630
 631 s_version 0x1234
 632 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 633
 634 s_waitcnt_expcnt exec_hi, 0x1234
 635 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 636
 637 s_waitcnt_lgkmcnt exec_hi, 0x1234
 638 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 639
 640 s_waitcnt_vmcnt exec_hi, 0x1234
 641 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 642
 643 s_waitcnt_vscnt exec_hi, 0x1234
 644 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 645
 646 s_xnor_saveexec_b32 exec_hi, s1
 647 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 648
 649 s_xor_saveexec_b32 exec_hi, s1
 650 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 651
 652 scratch_load_dword v0, v1, off offset:-2048 glc slc
 653 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 654
 655 scratch_load_dwordx2 v[1:2], v3, off
 656 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 657
 658 scratch_load_dwordx3 v[1:3], v4, off
 659 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 660
 661 scratch_load_dwordx4 v[1:4], v5, off
 662 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 663
 664 scratch_load_sbyte v1, v2, off
 665 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 666
 667 scratch_load_sbyte_d16 v1, v2, off
 668 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 669
 670 scratch_load_sbyte_d16_hi v1, v2, off
 671 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 672
 673 scratch_load_short_d16 v1, v2, off
 674 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 675
 676 scratch_load_short_d16_hi v1, v2, off
 677 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 678
 679 scratch_load_sshort v1, v2, off
 680 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 681
 682 scratch_load_ubyte v1, v2, off
 683 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 684
 685 scratch_load_ubyte_d16 v1, v2, off
 686 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 687
 688 scratch_load_ubyte_d16_hi v1, v2, off
 689 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 690
 691 scratch_load_ushort v1, v2, off
 692 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 693
 694 scratch_store_byte off, v2, flat_scratch_hi offset:-1
 695 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 696
 697 scratch_store_byte_d16_hi off, v2, flat_scratch_hi offset:-1
 698 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 699
 700 scratch_store_dword off, v2, exec_hi
 701 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 702
 703 scratch_store_dwordx2 off, v[254:255], s3 offset:-1
 704 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 705
 706 scratch_store_dwordx3 off, v[253:255], s3 offset:-1
 707 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 708
 709 scratch_store_dwordx4 off, v[252:255], s3 offset:-1
 710 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 711
 712 scratch_store_short off, v2, flat_scratch_hi offset:-1
 713 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 714
 715 scratch_store_short_d16_hi off, v2, flat_scratch_hi offset:-1
 716 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 717
 718 v_accvgpr_read_b32 a0, a0
 719 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 720
 721 v_accvgpr_write_b32 a0, 65
 722 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 723
 724 v_add3_u32 v1, v2, v3, v4
 725 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 726
 727 v_add_co_ci_u32 v1, sext(v1), sext(v4) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:DWORD
 728 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 729
 730 v_add_co_ci_u32_dpp v0, vcc, v0, v0, vcc dpp8:[7,6,5,4,3,2,1,0] fi:1
 731 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 732
 733 v_add_co_ci_u32_e32 v255, vcc, v1, v2, vcc
 734 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 735
 736 v_add_co_ci_u32_e64 v255, s12, v1, v2, s6
 737 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 738
 739 v_add_co_ci_u32_sdwa v1, v1, v4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:DWORD
 740 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 741
 742 v_add_co_u32 v0, exec, v0, v2
 743 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 744
 745 v_add_co_u32_dpp v255, vcc, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 746 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 747
 748 v_add_co_u32_e32 v2, vcc, s0, v2
 749 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 750
 751 v_add_co_u32_e64 v0, s0, v0, v2
 752 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 753
 754 v_add_co_u32_sdwa v0, v0, v4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:BYTE_0
 755 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 756
 757 v_add_i16 v255, v1, v2
 758 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 759
 760 v_add_i32 lds_direct, v0, v0
 761 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 762
 763 v_add_i32_e32 v0, vcc, 0.5, v0
 764 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 765
 766 v_add_i32_e64 v1, s[0:1], v2, v3
 767 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 768
 769 v_add_lshl_u32 v1, v2, v3, v4
 770 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 771
 772 v_add_nc_i16 v255, v1, v2
 773 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 774
 775 v_add_nc_i32 v255, v1, v2
 776 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 777
 778 v_add_nc_u16 v255, v1, v2
 779 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 780
 781 v_add_nc_u32_dpp v5, v1, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1
 782 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 783
 784 v_add_nc_u32_e32 v255, v1, v2
 785 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 786
 787 v_add_nc_u32_e64 v255, v1, v2
 788 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 789
 790 v_add_nc_u32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
 791 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 792
 793 v_addc_co_u32 v0, vcc, shared_base, v0, vcc
 794 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 795
 796 v_addc_co_u32_dpp v255, vcc, v1, v2, vcc quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 797 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 798
 799 v_addc_co_u32_e32 v3, vcc, 12345, v3, vcc
 800 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 801
 802 v_addc_co_u32_e64 v255, s[12:13], v1, v2, s[6:7]
 803 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 804
 805 v_addc_co_u32_sdwa v1, vcc, v2, v3, vcc dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
 806 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 807
 808 v_and_or_b32 v1, v2, v3, v4
 809 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 810
 811 v_ashr_i32 v255, v1, v2
 812 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 813
 814 v_ashr_i32_e64 v255, v1, v2
 815 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 816
 817 v_ashr_i64 v[254:255], v[1:2], v2
 818 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 819
 820 v_cmps_eq_f32 vcc, -1, v2
 821 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 822
 823 v_cmps_eq_f32_e64 flat_scratch, v1, v2
 824 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 825
 826 v_cmps_eq_f64 vcc, -1, v[2:3]
 827 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 828
 829 v_cmps_eq_f64_e64 flat_scratch, v[1:2], v[2:3]
 830 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 831
 832 v_cmps_f_f32 vcc, -1, v2
 833 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 834
 835 v_cmps_f_f32_e64 flat_scratch, v1, v2
 836 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 837
 838 v_cmps_f_f64 vcc, -1, v[2:3]
 839 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 840
 841 v_cmps_f_f64_e64 flat_scratch, v[1:2], v[2:3]
 842 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 843
 844 v_cmps_ge_f32 vcc, -1, v2
 845 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 846
 847 v_cmps_ge_f32_e64 flat_scratch, v1, v2
 848 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 849
 850 v_cmps_ge_f64 vcc, -1, v[2:3]
 851 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 852
 853 v_cmps_ge_f64_e64 flat_scratch, v[1:2], v[2:3]
 854 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 855
 856 v_cmps_gt_f32 vcc, -1, v2
 857 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 858
 859 v_cmps_gt_f32_e64 flat_scratch, v1, v2
 860 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 861
 862 v_cmps_gt_f64 vcc, -1, v[2:3]
 863 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 864
 865 v_cmps_gt_f64_e64 flat_scratch, v[1:2], v[2:3]
 866 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 867
 868 v_cmps_le_f32 vcc, -1, v2
 869 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 870
 871 v_cmps_le_f32_e64 flat_scratch, v1, v2
 872 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 873
 874 v_cmps_le_f64 vcc, -1, v[2:3]
 875 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 876
 877 v_cmps_le_f64_e64 flat_scratch, v[1:2], v[2:3]
 878 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 879
 880 v_cmps_lg_f32 vcc, -1, v2
 881 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 882
 883 v_cmps_lg_f32_e64 flat_scratch, v1, v2
 884 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 885
 886 v_cmps_lg_f64 vcc, -1, v[2:3]
 887 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 888
 889 v_cmps_lg_f64_e64 flat_scratch, v[1:2], v[2:3]
 890 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 891
 892 v_cmps_lt_f32 vcc, -1, v2
 893 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 894
 895 v_cmps_lt_f32_e64 flat_scratch, v1, v2
 896 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 897
 898 v_cmps_lt_f64 vcc, -1, v[2:3]
 899 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 900
 901 v_cmps_lt_f64_e64 flat_scratch, v[1:2], v[2:3]
 902 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 903
 904 v_cmps_neq_f32 vcc, -1, v2
 905 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 906
 907 v_cmps_neq_f32_e64 flat_scratch, v1, v2
 908 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 909
 910 v_cmps_neq_f64 vcc, -1, v[2:3]
 911 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 912
 913 v_cmps_neq_f64_e64 flat_scratch, v[1:2], v[2:3]
 914 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 915
 916 v_cmps_nge_f32 vcc, -1, v2
 917 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 918
 919 v_cmps_nge_f32_e64 flat_scratch, v1, v2
 920 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 921
 922 v_cmps_nge_f64 vcc, -1, v[2:3]
 923 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 924
 925 v_cmps_nge_f64_e64 flat_scratch, v[1:2], v[2:3]
 926 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 927
 928 v_cmps_ngt_f32 vcc, -1, v2
 929 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 930
 931 v_cmps_ngt_f32_e64 flat_scratch, v1, v2
 932 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 933
 934 v_cmps_ngt_f64 vcc, -1, v[2:3]
 935 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 936
 937 v_cmps_ngt_f64_e64 flat_scratch, v[1:2], v[2:3]
 938 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 939
 940 v_cmps_nle_f32 vcc, -1, v2
 941 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 942
 943 v_cmps_nle_f32_e64 flat_scratch, v1, v2
 944 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 945
 946 v_cmps_nle_f64 vcc, -1, v[2:3]
 947 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 948
 949 v_cmps_nle_f64_e64 flat_scratch, v[1:2], v[2:3]
 950 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 951
 952 v_cmps_nlg_f32 vcc, -1, v2
 953 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 954
 955 v_cmps_nlg_f32_e64 flat_scratch, v1, v2
 956 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 957
 958 v_cmps_nlg_f64 vcc, -1, v[2:3]
 959 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 960
 961 v_cmps_nlg_f64_e64 flat_scratch, v[1:2], v[2:3]
 962 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 963
 964 v_cmps_nlt_f32 vcc, -1, v2
 965 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 966
 967 v_cmps_nlt_f32_e64 flat_scratch, v1, v2
 968 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 969
 970 v_cmps_nlt_f64 vcc, -1, v[2:3]
 971 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 972
 973 v_cmps_nlt_f64_e64 flat_scratch, v[1:2], v[2:3]
 974 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 975
 976 v_cmps_o_f32 vcc, -1, v2
 977 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 978
 979 v_cmps_o_f32_e64 flat_scratch, v1, v2
 980 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 981
 982 v_cmps_o_f64 vcc, -1, v[2:3]
 983 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 984
 985 v_cmps_o_f64_e64 flat_scratch, v[1:2], v[2:3]
 986 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 987
 988 v_cmps_tru_f32 vcc, -1, v2
 989 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 990
 991 v_cmps_tru_f32_e64 flat_scratch, v1, v2
 992 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 993
 994 v_cmps_tru_f64 vcc, -1, v[2:3]
 995 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 996
 997 v_cmps_tru_f64_e64 flat_scratch, v[1:2], v[2:3]
 998 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 999
1000 v_cmps_u_f32 vcc, -1, v2
1001 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1002
1003 v_cmps_u_f32_e64 flat_scratch, v1, v2
1004 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1005
1006 v_cmps_u_f64 vcc, -1, v[2:3]
1007 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1008
1009 v_cmps_u_f64_e64 flat_scratch, v[1:2], v[2:3]
1010 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1011
1012 v_cmpsx_eq_f32 vcc, -1, v2
1013 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1014
1015 v_cmpsx_eq_f32_e64 flat_scratch, v1, v2
1016 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1017
1018 v_cmpsx_eq_f64 vcc, -1, v[2:3]
1019 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1020
1021 v_cmpsx_eq_f64_e64 flat_scratch, v[1:2], v[2:3]
1022 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1023
1024 v_cmpsx_f_f32 vcc, -1, v2
1025 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1026
1027 v_cmpsx_f_f32_e64 flat_scratch, v1, v2
1028 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1029
1030 v_cmpsx_f_f64 vcc, -1, v[2:3]
1031 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1032
1033 v_cmpsx_f_f64_e64 flat_scratch, v[1:2], v[2:3]
1034 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1035
1036 v_cmpsx_ge_f32 vcc, -1, v2
1037 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1038
1039 v_cmpsx_ge_f32_e64 flat_scratch, v1, v2
1040 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1041
1042 v_cmpsx_ge_f64 vcc, -1, v[2:3]
1043 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1044
1045 v_cmpsx_ge_f64_e64 flat_scratch, v[1:2], v[2:3]
1046 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1047
1048 v_cmpsx_gt_f32 vcc, -1, v2
1049 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1050
1051 v_cmpsx_gt_f32_e64 flat_scratch, v1, v2
1052 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1053
1054 v_cmpsx_gt_f64 vcc, -1, v[2:3]
1055 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1056
1057 v_cmpsx_gt_f64_e64 flat_scratch, v[1:2], v[2:3]
1058 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1059
1060 v_cmpsx_le_f32 vcc, -1, v2
1061 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1062
1063 v_cmpsx_le_f32_e64 flat_scratch, v1, v2
1064 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1065
1066 v_cmpsx_le_f64 vcc, -1, v[2:3]
1067 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1068
1069 v_cmpsx_le_f64_e64 flat_scratch, v[1:2], v[2:3]
1070 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1071
1072 v_cmpsx_lg_f32 vcc, -1, v2
1073 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1074
1075 v_cmpsx_lg_f32_e64 flat_scratch, v1, v2
1076 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1077
1078 v_cmpsx_lg_f64 vcc, -1, v[2:3]
1079 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1080
1081 v_cmpsx_lg_f64_e64 flat_scratch, v[1:2], v[2:3]
1082 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1083
1084 v_cmpsx_lt_f32 vcc, -1, v2
1085 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1086
1087 v_cmpsx_lt_f32_e64 flat_scratch, v1, v2
1088 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1089
1090 v_cmpsx_lt_f64 vcc, -1, v[2:3]
1091 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1092
1093 v_cmpsx_lt_f64_e64 flat_scratch, v[1:2], v[2:3]
1094 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1095
1096 v_cmpsx_neq_f32 vcc, -1, v2
1097 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1098
1099 v_cmpsx_neq_f32_e64 flat_scratch, v1, v2
1100 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1101
1102 v_cmpsx_neq_f64 vcc, -1, v[2:3]
1103 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1104
1105 v_cmpsx_neq_f64_e64 flat_scratch, v[1:2], v[2:3]
1106 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1107
1108 v_cmpsx_nge_f32 vcc, -1, v2
1109 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1110
1111 v_cmpsx_nge_f32_e64 flat_scratch, v1, v2
1112 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1113
1114 v_cmpsx_nge_f64 vcc, -1, v[2:3]
1115 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1116
1117 v_cmpsx_nge_f64_e64 flat_scratch, v[1:2], v[2:3]
1118 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1119
1120 v_cmpsx_ngt_f32 vcc, -1, v2
1121 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1122
1123 v_cmpsx_ngt_f32_e64 flat_scratch, v1, v2
1124 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1125
1126 v_cmpsx_ngt_f64 vcc, -1, v[2:3]
1127 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1128
1129 v_cmpsx_ngt_f64_e64 flat_scratch, v[1:2], v[2:3]
1130 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1131
1132 v_cmpsx_nle_f32 vcc, -1, v2
1133 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1134
1135 v_cmpsx_nle_f32_e64 flat_scratch, v1, v2
1136 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1137
1138 v_cmpsx_nle_f64 vcc, -1, v[2:3]
1139 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1140
1141 v_cmpsx_nle_f64_e64 flat_scratch, v[1:2], v[2:3]
1142 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1143
1144 v_cmpsx_nlg_f32 vcc, -1, v2
1145 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1146
1147 v_cmpsx_nlg_f32_e64 flat_scratch, v1, v2
1148 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1149
1150 v_cmpsx_nlg_f64 vcc, -1, v[2:3]
1151 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1152
1153 v_cmpsx_nlg_f64_e64 flat_scratch, v[1:2], v[2:3]
1154 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1155
1156 v_cmpsx_nlt_f32 vcc, -1, v2
1157 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1158
1159 v_cmpsx_nlt_f32_e64 flat_scratch, v1, v2
1160 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1161
1162 v_cmpsx_nlt_f64 vcc, -1, v[2:3]
1163 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1164
1165 v_cmpsx_nlt_f64_e64 flat_scratch, v[1:2], v[2:3]
1166 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1167
1168 v_cmpsx_o_f32 vcc, -1, v2
1169 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1170
1171 v_cmpsx_o_f32_e64 flat_scratch, v1, v2
1172 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1173
1174 v_cmpsx_o_f64 vcc, -1, v[2:3]
1175 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1176
1177 v_cmpsx_o_f64_e64 flat_scratch, v[1:2], v[2:3]
1178 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1179
1180 v_cmpsx_tru_f32 vcc, -1, v2
1181 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1182
1183 v_cmpsx_tru_f32_e64 flat_scratch, v1, v2
1184 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1185
1186 v_cmpsx_tru_f64 vcc, -1, v[2:3]
1187 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1188
1189 v_cmpsx_tru_f64_e64 flat_scratch, v[1:2], v[2:3]
1190 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1191
1192 v_cmpsx_u_f32 vcc, -1, v2
1193 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1194
1195 v_cmpsx_u_f32_e64 flat_scratch, v1, v2
1196 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1197
1198 v_cmpsx_u_f64 vcc, -1, v[2:3]
1199 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1200
1201 v_cmpsx_u_f64_e64 flat_scratch, v[1:2], v[2:3]
1202 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1203
1204 v_cvt_norm_i16_f16 v5, -4.0
1205 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1206
1207 v_cvt_norm_i16_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1208 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1209
1210 v_cvt_norm_i16_f16_e32 v255, v1
1211 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1212
1213 v_cvt_norm_i16_f16_e64 v255, v1
1214 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1215
1216 v_cvt_norm_i16_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1217 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1218
1219 v_cvt_norm_u16_f16 v5, s101
1220 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1221
1222 v_cvt_norm_u16_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1223 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1224
1225 v_cvt_norm_u16_f16_e32 v255, v1
1226 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1227
1228 v_cvt_norm_u16_f16_e64 v255, v1
1229 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1230
1231 v_cvt_norm_u16_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1232 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1233
1234 v_cvt_pknorm_i16_f16 v255, v1, v2
1235 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1236
1237 v_cvt_pknorm_u16_f16 v255, v1, v2
1238 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1239
1240 v_div_fixup_legacy_f16 v255, v1, v2, v3
1241 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1242
1243 v_div_fixup_legacy_f16_e64 v5, 0.5, v2, v3
1244 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1245
1246 v_dot2_f32_f16 v0, -v1, -v2, -v3
1247 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1248
1249 v_dot2_i32_i16 v0, -v1, -v2, -v3
1250 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1251
1252 v_dot2_u32_u16 v0, -v1, -v2, -v3
1253 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1254
1255 v_dot2c_f32_f16 v0, v1, v2
1256 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1257
1258 v_dot2c_f32_f16_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1259 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1260
1261 v_dot2c_f32_f16_e32 v255, v1, v2
1262 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1263
1264 v_dot2c_i32_i16 v0, v1, v2
1265 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1266
1267 v_dot2c_i32_i16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1268 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1269
1270 v_dot4_i32_i8 v0, v1, v2, v3
1271 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1272
1273 v_dot4_u32_u8 v0, v1, v2, v3
1274 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1275
1276 v_dot4c_i32_i8 v0, v1, v2
1277 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1278
1279 v_dot4c_i32_i8_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1280 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1281
1282 v_dot4c_i32_i8_e32 v255, v1, v2
1283 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1284
1285 v_dot8_i32_i4 v0, v1, v2, v3
1286 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1287
1288 v_dot8_u32_u4 v0, v1, v2, v3
1289 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1290
1291 v_dot8c_i32_i4 v0, v1, v2
1292 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1293
1294 v_dot8c_i32_i4_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1295 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1296
1297 v_fma_legacy_f16 v255, v1, v2, v3
1298 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1299
1300 v_fma_legacy_f16_e64 v5, v1, v2, v3
1301 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1302
1303 v_fma_mix_f32 v0, -abs(v1), v2, v3
1304 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1305
1306 v_fma_mixhi_f16 v0, -v1, abs(v2), -abs(v3)
1307 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1308
1309 v_fma_mixlo_f16 v0, abs(v1), -v2, abs(v3)
1310 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1311
1312 v_fmaak_f32 v255, v1, v2, 0x1121
1313 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1314
1315 v_fmac_f16 v5, 0x1234, v2
1316 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1317
1318 v_fmac_f16_dpp v5, v1, v2 quad_perm:[3,2,1,0] row_mask:0x0 bank_mask:0x0
1319 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1320
1321 v_fmac_f16_e32 v255, v1, v2
1322 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1323
1324 v_fmac_f16_e64 v255, v1, v2
1325 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1326
1327 v_fmac_f32 v0, v1, v2
1328 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1329
1330 v_fmac_f32_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1331 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1332
1333 v_fmac_f32_e32 v255, v1, v2
1334 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1335
1336 v_fmac_f32_e64 v255, v1, v2
1337 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1338
1339 v_fmamk_f32 v255, v1, 0x1121, v3
1340 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1341
1342 v_interp_p2_legacy_f16 v255, v2, attr0.x, v3
1343 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1344
1345 v_log_clamp_f32 v1, 0.5
1346 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1347
1348 v_log_clamp_f32_e64 v255, v1
1349 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1350
1351 v_lshl_add_u32 v1, v2, v3, v4
1352 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1353
1354 v_lshl_b32 v255, v1, v2
1355 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1356
1357 v_lshl_b32_e64 v255, v1, v2
1358 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1359
1360 v_lshl_b64 v[254:255], v[1:2], v2
1361 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1362
1363 v_lshl_or_b32 v1, v2, v3, v4
1364 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1365
1366 v_lshr_b32 v255, v1, v2
1367 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1368
1369 v_lshr_b32_e64 v255, v1, v2
1370 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1371
1372 v_lshr_b64 v[254:255], v[1:2], v2
1373 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1374
1375 v_mac_legacy_f32 v0, v1, v2
1376 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1377
1378 v_mac_legacy_f32_e32 v255, v1, v2
1379 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1380
1381 v_mac_legacy_f32_e64 v255, v1, v2
1382 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1383
1384 v_mad_i32_i16 v255, v1, v2, v3
1385 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1386
1387 v_mad_legacy_f16 v255, v1, v2, v3
1388 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1389
1390 v_mad_legacy_f16_e64 v5, 0.5, v2, v3
1391 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1392
1393 v_mad_legacy_i16 v255, v1, v2, v3
1394 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1395
1396 v_mad_legacy_i16_e64 v5, 0, v2, v3
1397 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1398
1399 v_mad_legacy_u16 v255, v1, v2, v3
1400 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1401
1402 v_mad_legacy_u16_e64 v5, 0, v2, v3
1403 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1404
1405 v_mad_mix_f32 v0, -abs(v1), v2, v3
1406 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1407
1408 v_mad_mixhi_f16 v0, -v1, abs(v2), -abs(v3)
1409 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1410
1411 v_mad_mixlo_f16 v0, abs(v1), -v2, abs(v3)
1412 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1413
1414 v_mad_u32_u16 v255, v1, v2, v3
1415 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1416
1417 v_max3_f16 v0, src_lds_direct, v0, v0
1418 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1419
1420 v_max3_i16 v1, v2, v3, v4
1421 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1422
1423 v_max3_u16 v1, v2, v3, v4
1424 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1425
1426 v_max_legacy_f32 v255, v1, v2
1427 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1428
1429 v_max_legacy_f32_e64 v255, v1, v2
1430 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1431
1432 v_med3_f16 v1, v2, v3, v4
1433 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1434
1435 v_med3_i16 v1, v2, v3, v4
1436 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1437
1438 v_med3_u16 v1, v2, v3, v4
1439 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1440
1441 v_mfma_f32_16x16x16f16 a[0:3], a[0:1], a[1:2], -2.0
1442 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1443
1444 v_mfma_f32_16x16x1f32 a[0:15], a0, a1, -2.0
1445 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1446
1447 v_mfma_f32_16x16x2bf16 a[0:15], a0, a1, -2.0
1448 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1449
1450 v_mfma_f32_16x16x4f16 a[0:15], a[0:1], a[1:2], -2.0
1451 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1452
1453 v_mfma_f32_16x16x4f32 a[0:3], a0, a1, -2.0
1454 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1455
1456 v_mfma_f32_16x16x8bf16 a[0:3], a0, a1, -2.0
1457 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1458
1459 v_mfma_f32_32x32x1f32 a[0:31], 1, v1, a[1:32]
1460 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1461
1462 v_mfma_f32_32x32x2bf16 a[0:31], a0, a1, -2.0
1463 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1464
1465 v_mfma_f32_32x32x2f32 a[0:15], a0, a1, -2.0
1466 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1467
1468 v_mfma_f32_32x32x4bf16 a[0:15], a0, a1, -2.0
1469 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1470
1471 v_mfma_f32_32x32x4f16 a[0:31], a[0:1], a[1:2], -2.0
1472 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1473
1474 v_mfma_f32_32x32x8f16 a[0:15], a[0:1], a[1:2], -2.0
1475 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1476
1477 v_mfma_f32_4x4x1f32 a[0:3], a0, a1, -2.0
1478 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1479
1480 v_mfma_f32_4x4x2bf16 a[0:3], a0, a1, -2.0
1481 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1482
1483 v_mfma_f32_4x4x4f16 a[0:3], a[0:1], a[1:2], -2.0
1484 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1485
1486 v_mfma_i32_16x16x16i8 a[0:3], a0, a1, 2
1487 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1488
1489 v_mfma_i32_16x16x4i8 a[0:15], a0, a1, 2
1490 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1491
1492 v_mfma_i32_32x32x4i8 a[0:31], a0, a1, 2
1493 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1494
1495 v_mfma_i32_32x32x8i8 a[0:15], a0, a1, 2
1496 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1497
1498 v_mfma_i32_4x4x4i8 a[0:3], a0, a1, 2
1499 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1500
1501 v_min3_f16 v1, v2, v3, v4
1502 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1503
1504 v_min3_i16 v0, src_lds_direct, v0, v0
1505 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1506
1507 v_min3_u16 v1, v2, v3, v4
1508 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1509
1510 v_min_legacy_f32 v255, v1, v2
1511 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1512
1513 v_min_legacy_f32_e64 v255, v1, v2
1514 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1515
1516 v_movrelsd_2_b32 v0, v255 dpp8:[7,6,5,4,3,2,1,0]
1517 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1518
1519 v_movrelsd_2_b32_dpp v0, v2 quad_perm:[3,2,1,0] row_mask:0x0 bank_mask:0x0
1520 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1521
1522 v_movrelsd_2_b32_e32 v5, 1
1523 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1524
1525 v_movrelsd_2_b32_e64 v255, v1
1526 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1527
1528 v_movrelsd_2_b32_sdwa v0, 0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1529 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1530
1531 v_mullit_f32 v255, v1, v2, v3
1532 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1533
1534 v_or3_b32 v1, v2, v3, v4
1535 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1536
1537 v_pack_b32_f16 v1, v2, v3
1538 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1539
1540 v_permlane16_b32 v0, lds_direct, s0, s0
1541 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1542
1543 v_permlanex16_b32 v0, lds_direct, s0, s0
1544 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1545
1546 v_pipeflush
1547 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1548
1549 v_pipeflush_e64
1550 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1551
1552 v_pk_add_f16 v0, execz, v0
1553 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1554
1555 v_pk_add_i16 v0, src_lds_direct, v0
1556 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1557
1558 v_pk_add_u16 v0, v1, v2
1559 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1560
1561 v_pk_ashrrev_i16 v0, lds_direct, v0
1562 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1563
1564 v_pk_fma_f16 v0, v1, v2, v3
1565 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1566
1567 v_pk_fmac_f16 v0, v1, v2
1568 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1569
1570 v_pk_lshlrev_b16 v0, lds_direct, v0
1571 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1572
1573 v_pk_lshrrev_b16 v0, lds_direct, v0
1574 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1575
1576 v_pk_mad_i16 v0, src_lds_direct, v0, v0
1577 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1578
1579 v_pk_mad_u16 v255, v1, v2, v3
1580 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1581
1582 v_pk_max_f16 v0, v1, v2
1583 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1584
1585 v_pk_max_i16 v0, v1, v2
1586 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1587
1588 v_pk_max_u16 v0, v1, v2
1589 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1590
1591 v_pk_min_f16 v0, v1, v2
1592 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1593
1594 v_pk_min_i16 v0, v1, v2
1595 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1596
1597 v_pk_min_u16 v0, v1, v2
1598 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1599
1600 v_pk_mul_f16 v0, v1, v2
1601 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1602
1603 v_pk_mul_lo_u16 v0, v1, v2
1604 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1605
1606 v_pk_sub_i16 v0, v1, v2
1607 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1608
1609 v_pk_sub_u16 v255, v1, v2
1610 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1611
1612 v_rcp_clamp_f32 v255, v1
1613 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1614
1615 v_rcp_clamp_f32_e64 v255, v1
1616 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1617
1618 v_rcp_clamp_f64 v[254:255], v[1:2]
1619 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1620
1621 v_rcp_clamp_f64_e64 v[254:255], v[1:2]
1622 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1623
1624 v_rcp_legacy_f32 v255, v1
1625 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1626
1627 v_rcp_legacy_f32_e64 v255, v1
1628 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1629
1630 v_rsq_clamp_f32 v255, v1
1631 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1632
1633 v_rsq_clamp_f32_e64 v255, v1
1634 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1635
1636 v_rsq_clamp_f64 v[254:255], v[1:2]
1637 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1638
1639 v_rsq_clamp_f64_e64 v[254:255], v[1:2]
1640 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1641
1642 v_rsq_legacy_f32 v255, v1
1643 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1644
1645 v_rsq_legacy_f32_e64 v255, v1
1646 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1647
1648 v_sat_pk_u8_i16 v255, v1
1649 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1650
1651 v_sat_pk_u8_i16_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0x0 bank_mask:0x0
1652 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1653
1654 v_sat_pk_u8_i16_e64 v5, -1
1655 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1656
1657 v_sat_pk_u8_i16_sdwa v5, sext(v1) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1658 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1659
1660 v_screen_partition_4se_b32 v5, -1
1661 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1662
1663 v_screen_partition_4se_b32_dpp v5, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0 bound_ctrl:0
1664 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1665
1666 v_screen_partition_4se_b32_e64 v5, -1
1667 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1668
1669 v_screen_partition_4se_b32_sdwa v5, v1 src0_sel:BYTE_0
1670 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1671
1672 v_sub_co_ci_u32_dpp v0, vcc, v0, v0, vcc dpp8:[7,6,5,4,3,2,1,0] fi:1
1673 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1674
1675 v_sub_co_ci_u32_e32 v255, vcc, v1, v2, vcc
1676 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1677
1678 v_sub_co_ci_u32_e64 v255, s12, v1, v2, s6
1679 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1680
1681 v_sub_co_ci_u32_sdwa v1, v1, v4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:DWORD
1682 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1683
1684 v_sub_co_u32 v0, s0, v0, v2
1685 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1686
1687 v_sub_co_u32_dpp v255, vcc, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1688 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1689
1690 v_sub_co_u32_e32 v2, vcc, s0, v2
1691 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1692
1693 v_sub_co_u32_e64 v0, s0, v0, v2
1694 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1695
1696 v_sub_co_u32_sdwa v0, v0, v4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:BYTE_0
1697 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1698
1699 v_sub_i16 v255, v1, v2
1700 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1701
1702 v_sub_i32 v1, s[0:1], v2, v3
1703 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1704
1705 v_sub_i32_e64 v255, s[12:13], v1, v2
1706 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1707
1708 v_sub_nc_i16 v255, v1, v2
1709 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1710
1711 v_sub_nc_i32 v255, v1, v2
1712 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1713
1714 v_sub_nc_u16 v255, v1, v2
1715 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1716
1717 v_sub_nc_u32_dpp v5, v1, v2 dpp8:[7,6,5,4,3,2,1,0]
1718 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1719
1720 v_sub_nc_u32_e32 v255, v1, v2
1721 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1722
1723 v_sub_nc_u32_e64 v255, v1, v2
1724 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1725
1726 v_sub_nc_u32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1727 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1728
1729 v_subb_co_u32 v1, vcc, v2, v3, vcc row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
1730 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1731
1732 v_subb_co_u32_dpp v255, vcc, v1, v2, vcc quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1733 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1734
1735 v_subb_co_u32_e64 v255, s[12:13], v1, v2, s[6:7]
1736 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1737
1738 v_subb_co_u32_sdwa v1, vcc, v2, v3, vcc dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
1739 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1740
1741 v_subbrev_co_u32 v0, vcc, src_lds_direct, v0, vcc
1742 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1743
1744 v_subbrev_co_u32_dpp v255, vcc, v1, v2, vcc quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1745 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1746
1747 v_subbrev_co_u32_e64 v255, s[12:13], v1, v2, s[6:7]
1748 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1749
1750 v_subbrev_co_u32_sdwa v1, vcc, v2, v3, vcc dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
1751 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1752
1753 v_subrev_co_ci_u32 v0, vcc_lo, src_lds_direct, v0, vcc_lo
1754 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1755
1756 v_subrev_co_ci_u32_dpp v0, vcc, v0, v0, vcc dpp8:[7,6,5,4,3,2,1,0]
1757 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1758
1759 v_subrev_co_ci_u32_e32 v1, 0, v1
1760 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1761
1762 v_subrev_co_ci_u32_e64 v255, s12, v1, v2, s6
1763 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1764
1765 v_subrev_co_ci_u32_sdwa v1, v1, v4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:DWORD
1766 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1767
1768 v_subrev_co_u32 v0, s0, src_lds_direct, v0
1769 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1770
1771 v_subrev_co_u32_dpp v255, vcc, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1772 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1773
1774 v_subrev_co_u32_e32 v2, vcc, s0, v2
1775 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1776
1777 v_subrev_co_u32_e64 v0, s0, v0, v2
1778 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1779
1780 v_subrev_co_u32_sdwa v0, v0, v4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:BYTE_0
1781 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1782
1783 v_subrev_i32 v1, s[0:1], v2, v3
1784 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1785
1786 v_subrev_i32_e64 v255, s[12:13], v1, v2
1787 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1788
1789 v_subrev_nc_u32 v0, src_lds_direct, v0
1790 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1791
1792 v_subrev_nc_u32_dpp v5, v1, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1
1793 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1794
1795 v_subrev_nc_u32_e32 v255, v1, v2
1796 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1797
1798 v_subrev_nc_u32_e64 v255, v1, v2
1799 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1800
1801 v_subrev_nc_u32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1802 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1803
1804 v_swap_b32 v1, 1
1805 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1806
1807 v_swap_b32_e32 v1, v2
1808 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1809
1810 v_swaprel_b32 v255, v1
1811 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1812
1813 v_xad_u32 v1, v2, v3, v4
1814 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1815
1816 v_xnor_b32 v0, v1, v2
1817 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1818
1819 v_xnor_b32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1820 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1821
1822 v_xnor_b32_e32 v255, v1, v2
1823 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1824
1825 v_xnor_b32_e64 v255, v1, v2
1826 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1827
1828 v_xnor_b32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1829 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1830
1831 v_xor3_b32 v255, v1, v2, v3
1832 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1833
1834 global_load_lds_dword v[2:3], off
1835 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1836
1837 global_load_dword v[2:3], off lds
1838 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1839
1840 scratch_load_dword v2, off lds
1841 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1842
1843 scratch_load_dword off, s2 lds
1844 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1845
1846 //===----------------------------------------------------------------------===//
1847 // Unsupported e32 variants.
1848 //===----------------------------------------------------------------------===//
1849
1850 v_cvt_pkrtz_f16_f32_e32 v255, v1, v2
1851 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: e32 variant of this instruction is not supported
1852
1853 //===----------------------------------------------------------------------===//
1854 // Unsupported dpp variants.
1855 //===----------------------------------------------------------------------===//
1856
1857 v_movreld_b32_dpp v1, v0 quad_perm:[3,2,1,0] row_mask:0x0 bank_mask:0x0
1858 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: dpp variant of this instruction is not supported
1859
1860 v_movrels_b32_dpp v1, v0 quad_perm:[3,2,1,0] row_mask:0x0 bank_mask:0x0 fi:1
1861 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: dpp variant of this instruction is not supported
1862
1863 v_movrelsd_b32_dpp v0, v255 quad_perm:[3,2,1,0] row_mask:0x0 bank_mask:0x0
1864 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: dpp variant of this instruction is not supported
1865
1866 //===----------------------------------------------------------------------===//
1867 // Unsupported sdwa variants.
1868 //===----------------------------------------------------------------------===//
1869
1870 v_movreld_b32_sdwa v0, 64 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1871 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: sdwa variant of this instruction is not supported
1872
1873 v_movrels_b32_sdwa v0, 1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1874 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: sdwa variant of this instruction is not supported
1875
1876 v_movrelsd_b32_sdwa v0, 1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1877 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: sdwa variant of this instruction is not supported