llvm/test/MC/AMDGPU/gfx7_unsupported.s

   1 // RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire %s 2>&1 | FileCheck --implicit-check-not=error: %s
   2
   3 //===----------------------------------------------------------------------===//
   4 // Unsupported instructions.
   5 //===----------------------------------------------------------------------===//
   6
   7 buffer_atomic_add_f32 v255, off, s[8:11], s3 offset:4095
   8 // CHECK: error: instruction not supported on this GPU
   9
  10 buffer_atomic_pk_add_f16 v255, off, s[8:11], s3 offset:4095
  11 // CHECK: error: instruction not supported on this GPU
  12
  13 buffer_gl0_inv
  14 // CHECK: error: instruction not supported on this GPU
  15
  16 buffer_gl1_inv
  17 // CHECK: error: instruction not supported on this GPU
  18
  19 buffer_load_format_d16_hi_x v5, off, s[8:11], s3
  20 // CHECK: error: instruction not supported on this GPU
  21
  22 buffer_load_format_d16_x v1, off, s[4:7], s1
  23 // CHECK: error: instruction not supported on this GPU
  24
  25 buffer_load_format_d16_xy v1, off, s[4:7], s1
  26 // CHECK: error: instruction not supported on this GPU
  27
  28 buffer_load_format_d16_xyz v[1:2], off, s[4:7], s1
  29 // CHECK: error: instruction not supported on this GPU
  30
  31 buffer_load_format_d16_xyzw v[1:2], off, s[4:7], s1
  32 // CHECK: error: instruction not supported on this GPU
  33
  34 buffer_load_sbyte_d16 v1, off, s[4:7], s1
  35 // CHECK: error: instruction not supported on this GPU
  36
  37 buffer_load_sbyte_d16_hi v1, off, s[4:7], s1
  38 // CHECK: error: instruction not supported on this GPU
  39
  40 buffer_load_short_d16 v1, off, s[4:7], s1
  41 // CHECK: error: instruction not supported on this GPU
  42
  43 buffer_load_short_d16_hi v1, off, s[4:7], s1
  44 // CHECK: error: instruction not supported on this GPU
  45
  46 buffer_load_ubyte_d16 v1, off, s[4:7], s1
  47 // CHECK: error: instruction not supported on this GPU
  48
  49 buffer_load_ubyte_d16_hi v1, off, s[4:7], s1
  50 // CHECK: error: instruction not supported on this GPU
  51
  52 buffer_store_byte_d16_hi v1, off, s[12:15], -1 offset:4095
  53 // CHECK: error: instruction not supported on this GPU
  54
  55 buffer_store_format_d16_hi_x v1, off, s[12:15], s4 offset:4095 glc
  56 // CHECK: error: instruction not supported on this GPU
  57
  58 buffer_store_format_d16_x v1, off, s[12:15], -1 offset:4095
  59 // CHECK: error: instruction not supported on this GPU
  60
  61 buffer_store_format_d16_xy v1, off, s[12:15], -1 offset:4095
  62 // CHECK: error: instruction not supported on this GPU
  63
  64 buffer_store_format_d16_xyz v[1:2], off, s[12:15], -1 offset:4095
  65 // CHECK: error: instruction not supported on this GPU
  66
  67 buffer_store_format_d16_xyzw v[1:2], off, s[12:15], -1 offset:4095
  68 // CHECK: error: instruction not supported on this GPU
  69
  70 buffer_store_lds_dword s[4:7], s0 lds
  71 // CHECK: error: instruction not supported on this GPU
  72
  73 buffer_store_short_d16_hi v1, off, s[12:15], -1 offset:4095
  74 // CHECK: error: instruction not supported on this GPU
  75
  76 ds_add_f32 v0, v1
  77 // CHECK: error: instruction not supported on this GPU
  78
  79 ds_add_rtn_f32 v0, v1, v2
  80 // CHECK: error: instruction not supported on this GPU
  81
  82 ds_add_src2_f32 v0 offset:4 gds
  83 // CHECK: error: instruction not supported on this GPU
  84
  85 ds_bpermute_b32 v0, v1, v2
  86 // CHECK: error: instruction not supported on this GPU
  87
  88 ds_permute_b32 v0, v1, v2
  89 // CHECK: error: instruction not supported on this GPU
  90
  91 ds_read_addtid_b32 v255 offset:65535
  92 // CHECK: error: instruction not supported on this GPU
  93
  94 ds_read_i8_d16 v255, v1 offset:65535
  95 // CHECK: error: instruction not supported on this GPU
  96
  97 ds_read_i8_d16_hi v255, v1 offset:65535
  98 // CHECK: error: instruction not supported on this GPU
  99
 100 ds_read_u16_d16 v255, v1 offset:65535
 101 // CHECK: error: instruction not supported on this GPU
 102
 103 ds_read_u16_d16_hi v255, v1 offset:65535
 104 // CHECK: error: instruction not supported on this GPU
 105
 106 ds_read_u8_d16 v255, v1 offset:65535
 107 // CHECK: error: instruction not supported on this GPU
 108
 109 ds_read_u8_d16_hi v255, v1 offset:65535
 110 // CHECK: error: instruction not supported on this GPU
 111
 112 ds_write_addtid_b32 v255 offset:65535
 113 // CHECK: error: instruction not supported on this GPU
 114
 115 ds_write_b16_d16_hi v1, v2
 116 // CHECK: error: instruction not supported on this GPU
 117
 118 ds_write_b8_d16_hi v1, v2
 119 // CHECK: error: instruction not supported on this GPU
 120
 121 flat_load_sbyte_d16 v1, v[3:4]
 122 // CHECK: error: instruction not supported on this GPU
 123
 124 flat_load_sbyte_d16_hi v1, v[3:4]
 125 // CHECK: error: instruction not supported on this GPU
 126
 127 flat_load_short_d16 v1, v[3:4]
 128 // CHECK: error: instruction not supported on this GPU
 129
 130 flat_load_short_d16_hi v1, v[3:4]
 131 // CHECK: error: instruction not supported on this GPU
 132
 133 flat_load_ubyte_d16 v1, v[3:4]
 134 // CHECK: error: instruction not supported on this GPU
 135
 136 flat_load_ubyte_d16_hi v1, v[3:4]
 137 // CHECK: error: instruction not supported on this GPU
 138
 139 flat_store_byte_d16_hi v[1:2], v2
 140 // CHECK: error: instruction not supported on this GPU
 141
 142 flat_store_short_d16_hi v[1:2], v2
 143 // CHECK: error: instruction not supported on this GPU
 144
 145 global_atomic_add v0, v[1:2], v2, off glc slc
 146 // CHECK: error: instruction not supported on this GPU
 147
 148 global_atomic_add_f32 v[1:2], v2, off
 149 // CHECK: error: instruction not supported on this GPU
 150
 151 global_atomic_add_x2 v[1:2], v[254:255], off offset:-1
 152 // CHECK: error: instruction not supported on this GPU
 153
 154 global_atomic_and v[1:2], v2, off
 155 // CHECK: error: instruction not supported on this GPU
 156
 157 global_atomic_and_x2 v[1:2], v[254:255], off offset:-1
 158 // CHECK: error: instruction not supported on this GPU
 159
 160 global_atomic_cmpswap v[1:2], v[254:255], off offset:-1
 161 // CHECK: error: instruction not supported on this GPU
 162
 163 global_atomic_cmpswap_x2 v[1:2], v[252:255], off offset:-1
 164 // CHECK: error: instruction not supported on this GPU
 165
 166 global_atomic_dec v[1:2], v2, off
 167 // CHECK: error: instruction not supported on this GPU
 168
 169 global_atomic_dec_x2 v[1:2], v[254:255], off offset:-1
 170 // CHECK: error: instruction not supported on this GPU
 171
 172 global_atomic_inc v[1:2], v2, off
 173 // CHECK: error: instruction not supported on this GPU
 174
 175 global_atomic_inc_x2 v[1:2], v[254:255], off offset:-1
 176 // CHECK: error: instruction not supported on this GPU
 177
 178 global_atomic_or v[1:2], v2, off
 179 // CHECK: error: instruction not supported on this GPU
 180
 181 global_atomic_or_x2 v[1:2], v[254:255], off offset:-1
 182 // CHECK: error: instruction not supported on this GPU
 183
 184 global_atomic_pk_add_f16 v[1:2], v2, off
 185 // CHECK: error: instruction not supported on this GPU
 186
 187 global_atomic_smax v[1:2], v2, off
 188 // CHECK: error: instruction not supported on this GPU
 189
 190 global_atomic_smax_x2 v[1:2], v[254:255], off offset:-1
 191 // CHECK: error: instruction not supported on this GPU
 192
 193 global_atomic_smin v[1:2], v2, off
 194 // CHECK: error: instruction not supported on this GPU
 195
 196 global_atomic_smin_x2 v[1:2], v[254:255], off offset:-1
 197 // CHECK: error: instruction not supported on this GPU
 198
 199 global_atomic_sub v[1:2], v2, off
 200 // CHECK: error: instruction not supported on this GPU
 201
 202 global_atomic_sub_x2 v[1:2], v[254:255], off offset:-1
 203 // CHECK: error: instruction not supported on this GPU
 204
 205 global_atomic_swap v[1:2], v2, off
 206 // CHECK: error: instruction not supported on this GPU
 207
 208 global_atomic_swap_x2 v[1:2], v[254:255], off offset:-1
 209 // CHECK: error: instruction not supported on this GPU
 210
 211 global_atomic_umax v[1:2], v2, off
 212 // CHECK: error: instruction not supported on this GPU
 213
 214 global_atomic_umax_x2 v[1:2], v[254:255], off offset:-1
 215 // CHECK: error: instruction not supported on this GPU
 216
 217 global_atomic_umin v[1:2], v2, off
 218 // CHECK: error: instruction not supported on this GPU
 219
 220 global_atomic_umin_x2 v[1:2], v[254:255], off offset:-1
 221 // CHECK: error: instruction not supported on this GPU
 222
 223 global_atomic_xor v[1:2], v2, off
 224 // CHECK: error: instruction not supported on this GPU
 225
 226 global_atomic_xor_x2 v[1:2], v[254:255], off offset:-1
 227 // CHECK: error: instruction not supported on this GPU
 228
 229 global_load_dword v1, v3, s[2:3]
 230 // CHECK: error: instruction not supported on this GPU
 231
 232 global_load_dwordx2 v[1:2], v[3:4], off
 233 // CHECK: error: instruction not supported on this GPU
 234
 235 global_load_dwordx3 v[1:3], v[3:4], off
 236 // CHECK: error: instruction not supported on this GPU
 237
 238 global_load_dwordx4 v[1:4], v[3:4], off
 239 // CHECK: error: instruction not supported on this GPU
 240
 241 global_load_sbyte v1, v[3:4], off
 242 // CHECK: error: instruction not supported on this GPU
 243
 244 global_load_sbyte_d16 v1, v[3:4], off
 245 // CHECK: error: instruction not supported on this GPU
 246
 247 global_load_sbyte_d16_hi v1, v[3:4], off
 248 // CHECK: error: instruction not supported on this GPU
 249
 250 global_load_short_d16 v1, v[3:4], off
 251 // CHECK: error: instruction not supported on this GPU
 252
 253 global_load_short_d16_hi v1, v[3:4], off
 254 // CHECK: error: instruction not supported on this GPU
 255
 256 global_load_sshort v1, v[3:4], off
 257 // CHECK: error: instruction not supported on this GPU
 258
 259 global_load_ubyte v1, v[3:4], off
 260 // CHECK: error: instruction not supported on this GPU
 261
 262 global_load_ubyte_d16 v1, v[3:4], off
 263 // CHECK: error: instruction not supported on this GPU
 264
 265 global_load_ubyte_d16_hi v1, v[3:4], off
 266 // CHECK: error: instruction not supported on this GPU
 267
 268 global_load_ushort v1, v[3:4], off
 269 // CHECK: error: instruction not supported on this GPU
 270
 271 global_store_byte v[1:2], v2, off
 272 // CHECK: error: instruction not supported on this GPU
 273
 274 global_store_byte_d16_hi v[1:2], v2, off
 275 // CHECK: error: instruction not supported on this GPU
 276
 277 global_store_dword v254, v1, s[2:3] offset:16
 278 // CHECK: error: instruction not supported on this GPU
 279
 280 global_store_dwordx2 v[1:2], v[254:255], off offset:-1
 281 // CHECK: error: instruction not supported on this GPU
 282
 283 global_store_dwordx3 v[1:2], v[253:255], off offset:-1
 284 // CHECK: error: instruction not supported on this GPU
 285
 286 global_store_dwordx4 v[1:2], v[252:255], off offset:-1
 287 // CHECK: error: instruction not supported on this GPU
 288
 289 global_store_short v[1:2], v2, off
 290 // CHECK: error: instruction not supported on this GPU
 291
 292 global_store_short_d16_hi v[1:2], v2, off
 293 // CHECK: error: instruction not supported on this GPU
 294
 295 s_and_saveexec_b32 exec_hi, s1
 296 // CHECK: error: instruction not supported on this GPU
 297
 298 s_andn1_saveexec_b32 exec_hi, s1
 299 // CHECK: error: instruction not supported on this GPU
 300
 301 s_andn1_saveexec_b64 exec, s[2:3]
 302 // CHECK: error: instruction not supported on this GPU
 303
 304 s_andn1_wrexec_b32 exec_hi, s1
 305 // CHECK: error: instruction not supported on this GPU
 306
 307 s_andn1_wrexec_b64 exec, s[2:3]
 308 // CHECK: error: instruction not supported on this GPU
 309
 310 s_andn2_saveexec_b32 exec_hi, s1
 311 // CHECK: error: instruction not supported on this GPU
 312
 313 s_andn2_wrexec_b32 exec_hi, s1
 314 // CHECK: error: instruction not supported on this GPU
 315
 316 s_andn2_wrexec_b64 exec, s[2:3]
 317 // CHECK: error: instruction not supported on this GPU
 318
 319 s_atc_probe 0x0, s[4:5], 0x0
 320 // CHECK: error: instruction not supported on this GPU
 321
 322 s_atc_probe_buffer 0x0, s[8:11], s101
 323 // CHECK: error: instruction not supported on this GPU
 324
 325 s_atomic_add s5, s[2:3], 0x0
 326 // CHECK: error: instruction not supported on this GPU
 327
 328 s_atomic_add_x2 s[10:11], s[2:3], s101
 329 // CHECK: error: instruction not supported on this GPU
 330
 331 s_atomic_and s5, s[2:3], s101
 332 // CHECK: error: instruction not supported on this GPU
 333
 334 s_atomic_and_x2 s[10:11], s[2:3], 0x0
 335 // CHECK: error: instruction not supported on this GPU
 336
 337 s_atomic_cmpswap s[10:11], s[2:3], 0x0
 338 // CHECK: error: instruction not supported on this GPU
 339
 340 s_atomic_cmpswap_x2 s[20:23], s[2:3], 0x0
 341 // CHECK: error: instruction not supported on this GPU
 342
 343 s_atomic_dec s5, s[2:3], s0 glc
 344 // CHECK: error: instruction not supported on this GPU
 345
 346 s_atomic_dec_x2 s[10:11], s[2:3], s101
 347 // CHECK: error: instruction not supported on this GPU
 348
 349 s_atomic_inc s5, s[2:3], s0 glc
 350 // CHECK: error: instruction not supported on this GPU
 351
 352 s_atomic_inc_x2 s[10:11], s[2:3], s101
 353 // CHECK: error: instruction not supported on this GPU
 354
 355 s_atomic_or s5, s[2:3], 0x0
 356 // CHECK: error: instruction not supported on this GPU
 357
 358 s_atomic_or_x2 s[10:11], s[2:3], s0 glc
 359 // CHECK: error: instruction not supported on this GPU
 360
 361 s_atomic_smax s5, s[2:3], s101
 362 // CHECK: error: instruction not supported on this GPU
 363
 364 s_atomic_smax_x2 s[10:11], s[2:3], s0 glc
 365 // CHECK: error: instruction not supported on this GPU
 366
 367 s_atomic_smin s5, s[2:3], s101
 368 // CHECK: error: instruction not supported on this GPU
 369
 370 s_atomic_smin_x2 s[10:11], s[2:3], s0 glc
 371 // CHECK: error: instruction not supported on this GPU
 372
 373 s_atomic_sub s5, s[2:3], s101
 374 // CHECK: error: instruction not supported on this GPU
 375
 376 s_atomic_sub_x2 s[10:11], s[2:3], s0 glc
 377 // CHECK: error: instruction not supported on this GPU
 378
 379 s_atomic_swap s5, s[2:3], -1
 380 // CHECK: error: instruction not supported on this GPU
 381
 382 s_atomic_swap_x2 s[10:11], s[2:3], s0 glc
 383 // CHECK: error: instruction not supported on this GPU
 384
 385 s_atomic_umax s5, s[2:3], s0 glc
 386 // CHECK: error: instruction not supported on this GPU
 387
 388 s_atomic_umax_x2 s[10:11], s[2:3], s101
 389 // CHECK: error: instruction not supported on this GPU
 390
 391 s_atomic_umin s5, s[2:3], s101
 392 // CHECK: error: instruction not supported on this GPU
 393
 394 s_atomic_umin_x2 s[10:11], s[2:3], s0 glc
 395 // CHECK: error: instruction not supported on this GPU
 396
 397 s_atomic_xor s5, s[2:3], s101
 398 // CHECK: error: instruction not supported on this GPU
 399
 400 s_atomic_xor_x2 s[10:11], s[2:3], s0 glc
 401 // CHECK: error: instruction not supported on this GPU
 402
 403 s_bitreplicate_b64_b32 exec, s2
 404 // CHECK: error: instruction not supported on this GPU
 405
 406 s_buffer_atomic_add s5, s[4:7], 0x0
 407 // CHECK: error: instruction not supported on this GPU
 408
 409 s_buffer_atomic_add_x2 s[10:11], s[4:7], s0
 410 // CHECK: error: instruction not supported on this GPU
 411
 412 s_buffer_atomic_and s101, s[4:7], s0
 413 // CHECK: error: instruction not supported on this GPU
 414
 415 s_buffer_atomic_and_x2 s[10:11], s[8:11], s0
 416 // CHECK: error: instruction not supported on this GPU
 417
 418 s_buffer_atomic_cmpswap s[10:11], s[4:7], 0x0
 419 // CHECK: error: instruction not supported on this GPU
 420
 421 s_buffer_atomic_cmpswap_x2 s[20:23], s[4:7], 0x0
 422 // CHECK: error: instruction not supported on this GPU
 423
 424 s_buffer_atomic_dec s5, s[4:7], s0
 425 // CHECK: error: instruction not supported on this GPU
 426
 427 s_buffer_atomic_dec_x2 s[10:11], s[4:7], s0 glc
 428 // CHECK: error: instruction not supported on this GPU
 429
 430 s_buffer_atomic_inc s101, s[4:7], s0
 431 // CHECK: error: instruction not supported on this GPU
 432
 433 s_buffer_atomic_inc_x2 s[10:11], s[4:7], 0x0
 434 // CHECK: error: instruction not supported on this GPU
 435
 436 s_buffer_atomic_or s5, s[8:11], s0
 437 // CHECK: error: instruction not supported on this GPU
 438
 439 s_buffer_atomic_or_x2 s[10:11], s[96:99], s0
 440 // CHECK: error: instruction not supported on this GPU
 441
 442 s_buffer_atomic_smax s5, s[4:7], s101
 443 // CHECK: error: instruction not supported on this GPU
 444
 445 s_buffer_atomic_smax_x2 s[100:101], s[4:7], s0
 446 // CHECK: error: instruction not supported on this GPU
 447
 448 s_buffer_atomic_smin s5, s[4:7], 0x0
 449 // CHECK: error: instruction not supported on this GPU
 450
 451 s_buffer_atomic_smin_x2 s[12:13], s[4:7], s0
 452 // CHECK: error: instruction not supported on this GPU
 453
 454 s_buffer_atomic_sub s5, s[4:7], s0 glc
 455 // CHECK: error: instruction not supported on this GPU
 456
 457 s_buffer_atomic_sub_x2 s[10:11], s[4:7], s0
 458 // CHECK: error: instruction not supported on this GPU
 459
 460 s_buffer_atomic_swap s5, s[4:7], -1
 461 // CHECK: error: instruction not supported on this GPU
 462
 463 s_buffer_atomic_swap_x2 s[10:11], s[4:7], s0 glc
 464 // CHECK: error: instruction not supported on this GPU
 465
 466 s_buffer_atomic_umax s5, s[4:7], s0
 467 // CHECK: error: instruction not supported on this GPU
 468
 469 s_buffer_atomic_umax_x2 s[10:11], s[4:7], s0 glc
 470 // CHECK: error: instruction not supported on this GPU
 471
 472 s_buffer_atomic_umin s5, s[4:7], s0
 473 // CHECK: error: instruction not supported on this GPU
 474
 475 s_buffer_atomic_umin_x2 s[10:11], s[4:7], s0 glc
 476 // CHECK: error: instruction not supported on this GPU
 477
 478 s_buffer_atomic_xor s5, s[4:7], s0
 479 // CHECK: error: instruction not supported on this GPU
 480
 481 s_buffer_atomic_xor_x2 s[10:11], s[4:7], s0 glc
 482 // CHECK: error: instruction not supported on this GPU
 483
 484 s_buffer_store_dword exec_hi, s[0:3], 0x0
 485 // CHECK: error: instruction not supported on this GPU
 486
 487 s_buffer_store_dwordx2 exec, s[0:3], 0x0
 488 // CHECK: error: instruction not supported on this GPU
 489
 490 s_buffer_store_dwordx4 s[4:7], s[12:15], m0
 491 // CHECK: error: instruction not supported on this GPU
 492
 493 s_call_b64 exec, 0x1234
 494 // CHECK: error: instruction not supported on this GPU
 495
 496 s_clause 0x0
 497 // CHECK: error: instruction not supported on this GPU
 498
 499 s_cmp_eq_u64 -1, s[4:5]
 500 // CHECK: error: instruction not supported on this GPU
 501
 502 s_cmp_lg_u64 -1, s[4:5]
 503 // CHECK: error: instruction not supported on this GPU
 504
 505 s_code_end
 506 // CHECK: error: instruction not supported on this GPU
 507
 508 s_dcache_discard s[2:3], 0x0
 509 // CHECK: error: instruction not supported on this GPU
 510
 511 s_dcache_discard_x2 s[2:3], 0x0
 512 // CHECK: error: instruction not supported on this GPU
 513
 514 s_dcache_wb
 515 // CHECK: error: instruction not supported on this GPU
 516
 517 s_dcache_wb_vol
 518 // CHECK: error: instruction not supported on this GPU
 519
 520 s_denorm_mode 0x0
 521 // CHECK: error: instruction not supported on this GPU
 522
 523 s_endpgm_ordered_ps_done
 524 // CHECK: error: instruction not supported on this GPU
 525
 526 s_endpgm_saved
 527 // CHECK: error: instruction not supported on this GPU
 528
 529 s_get_waveid_in_workgroup s0
 530 // CHECK: error: instruction not supported on this GPU
 531
 532 s_gl1_inv
 533 // CHECK: error: instruction not supported on this GPU
 534
 535 s_inst_prefetch 0x0
 536 // CHECK: error: instruction not supported on this GPU
 537
 538 s_lshl1_add_u32 exec_hi, s1, s2
 539 // CHECK: error: instruction not supported on this GPU
 540
 541 s_lshl2_add_u32 exec_hi, s1, s2
 542 // CHECK: error: instruction not supported on this GPU
 543
 544 s_lshl3_add_u32 exec_hi, s1, s2
 545 // CHECK: error: instruction not supported on this GPU
 546
 547 s_lshl4_add_u32 exec_hi, s1, s2
 548 // CHECK: error: instruction not supported on this GPU
 549
 550 s_memrealtime exec
 551 // CHECK: error: instruction not supported on this GPU
 552
 553 s_movrelsd_2_b32 s0, s1
 554 // CHECK: error: instruction not supported on this GPU
 555
 556 s_mul_hi_i32 exec_hi, s1, s2
 557 // CHECK: error: instruction not supported on this GPU
 558
 559 s_mul_hi_u32 exec_hi, s1, s2
 560 // CHECK: error: instruction not supported on this GPU
 561
 562 s_nand_saveexec_b32 exec_hi, s1
 563 // CHECK: error: instruction not supported on this GPU
 564
 565 s_nor_saveexec_b32 exec_hi, s1
 566 // CHECK: error: instruction not supported on this GPU
 567
 568 s_or_saveexec_b32 exec_hi, s1
 569 // CHECK: error: instruction not supported on this GPU
 570
 571 s_orn1_saveexec_b32 exec_hi, s1
 572 // CHECK: error: instruction not supported on this GPU
 573
 574 s_orn1_saveexec_b64 exec, s[2:3]
 575 // CHECK: error: instruction not supported on this GPU
 576
 577 s_orn2_saveexec_b32 exec_hi, s1
 578 // CHECK: error: instruction not supported on this GPU
 579
 580 s_pack_hh_b32_b16 exec_hi, s1, s2
 581 // CHECK: error: instruction not supported on this GPU
 582
 583 s_pack_lh_b32_b16 exec_hi, s1, s2
 584 // CHECK: error: instruction not supported on this GPU
 585
 586 s_pack_ll_b32_b16 exec_hi, s1, s2
 587 // CHECK: error: instruction not supported on this GPU
 588
 589 s_rfe_restore_b64 -1, s2
 590 // CHECK: error: instruction not supported on this GPU
 591
 592 s_round_mode 0x0
 593 // CHECK: error: instruction not supported on this GPU
 594
 595 s_scratch_load_dword s5, s[2:3], s0 glc
 596 // CHECK: error: instruction not supported on this GPU
 597
 598 s_scratch_load_dwordx2 s[100:101], s[2:3], s0
 599 // CHECK: error: instruction not supported on this GPU
 600
 601 s_scratch_load_dwordx4 s[20:23], s[4:5], s0
 602 // CHECK: error: instruction not supported on this GPU
 603
 604 s_scratch_store_dword s1, s[4:5], 0x123 glc
 605 // CHECK: error: instruction not supported on this GPU
 606
 607 s_scratch_store_dwordx2 s[2:3], s[4:5], s101 glc
 608 // CHECK: error: instruction not supported on this GPU
 609
 610 s_scratch_store_dwordx4 s[4:7], s[4:5], s0 glc
 611 // CHECK: error: instruction not supported on this GPU
 612
 613 s_set_gpr_idx_idx -1
 614 // CHECK: error: instruction not supported on this GPU
 615
 616 s_set_gpr_idx_mode 0
 617 // CHECK: error: instruction not supported on this GPU
 618
 619 s_set_gpr_idx_off
 620 // CHECK: error: instruction not supported on this GPU
 621
 622 s_set_gpr_idx_on -1, 0x0
 623 // CHECK: error: instruction not supported on this GPU
 624
 625 s_store_dword exec_hi, s[2:3], 0x0
 626 // CHECK: error: instruction not supported on this GPU
 627
 628 s_store_dwordx2 exec, s[2:3], 0x0
 629 // CHECK: error: instruction not supported on this GPU
 630
 631 s_store_dwordx4 s[4:7], flat_scratch, m0
 632 // CHECK: error: instruction not supported on this GPU
 633
 634 s_subvector_loop_begin exec_hi, 0x1234
 635 // CHECK: error: instruction not supported on this GPU
 636
 637 s_subvector_loop_end exec_hi, 0x1234
 638 // CHECK: error: instruction not supported on this GPU
 639
 640 s_ttracedata_imm 0x0
 641 // CHECK: error: instruction not supported on this GPU
 642
 643 s_version 0x1234
 644 // CHECK: error: instruction not supported on this GPU
 645
 646 s_waitcnt_expcnt exec_hi, 0x1234
 647 // CHECK: error: instruction not supported on this GPU
 648
 649 s_waitcnt_lgkmcnt exec_hi, 0x1234
 650 // CHECK: error: instruction not supported on this GPU
 651
 652 s_waitcnt_vmcnt exec_hi, 0x1234
 653 // CHECK: error: instruction not supported on this GPU
 654
 655 s_waitcnt_vscnt exec_hi, 0x1234
 656 // CHECK: error: instruction not supported on this GPU
 657
 658 s_wakeup
 659 // CHECK: error: instruction not supported on this GPU
 660
 661 s_xnor_saveexec_b32 exec_hi, s1
 662 // CHECK: error: instruction not supported on this GPU
 663
 664 s_xor_saveexec_b32 exec_hi, s1
 665 // CHECK: error: instruction not supported on this GPU
 666
 667 scratch_load_dword v0, v1, off offset:-2048 glc slc
 668 // CHECK: error: instruction not supported on this GPU
 669
 670 scratch_load_dwordx2 v[1:2], v3, off
 671 // CHECK: error: instruction not supported on this GPU
 672
 673 scratch_load_dwordx3 v[1:3], v4, off
 674 // CHECK: error: instruction not supported on this GPU
 675
 676 scratch_load_dwordx4 v[1:4], v5, off
 677 // CHECK: error: instruction not supported on this GPU
 678
 679 scratch_load_sbyte v1, v2, off
 680 // CHECK: error: instruction not supported on this GPU
 681
 682 scratch_load_sbyte_d16 v1, v2, off
 683 // CHECK: error: instruction not supported on this GPU
 684
 685 scratch_load_sbyte_d16_hi v1, v2, off
 686 // CHECK: error: instruction not supported on this GPU
 687
 688 scratch_load_short_d16 v1, v2, off
 689 // CHECK: error: instruction not supported on this GPU
 690
 691 scratch_load_short_d16_hi v1, v2, off
 692 // CHECK: error: instruction not supported on this GPU
 693
 694 scratch_load_sshort v1, v2, off
 695 // CHECK: error: instruction not supported on this GPU
 696
 697 scratch_load_ubyte v1, v2, off
 698 // CHECK: error: instruction not supported on this GPU
 699
 700 scratch_load_ubyte_d16 v1, v2, off
 701 // CHECK: error: instruction not supported on this GPU
 702
 703 scratch_load_ubyte_d16_hi v1, v2, off
 704 // CHECK: error: instruction not supported on this GPU
 705
 706 scratch_load_ushort v1, v2, off
 707 // CHECK: error: instruction not supported on this GPU
 708
 709 scratch_store_byte off, v2, flat_scratch_hi offset:-1
 710 // CHECK: error: instruction not supported on this GPU
 711
 712 scratch_store_byte_d16_hi off, v2, flat_scratch_hi offset:-1
 713 // CHECK: error: instruction not supported on this GPU
 714
 715 scratch_store_dword off, v2, exec_hi
 716 // CHECK: error: instruction not supported on this GPU
 717
 718 scratch_store_dwordx2 off, v[254:255], s3 offset:-1
 719 // CHECK: error: instruction not supported on this GPU
 720
 721 scratch_store_dwordx3 off, v[253:255], s3 offset:-1
 722 // CHECK: error: instruction not supported on this GPU
 723
 724 scratch_store_dwordx4 off, v[252:255], s3 offset:-1
 725 // CHECK: error: instruction not supported on this GPU
 726
 727 scratch_store_short off, v2, flat_scratch_hi offset:-1
 728 // CHECK: error: instruction not supported on this GPU
 729
 730 scratch_store_short_d16_hi off, v2, flat_scratch_hi offset:-1
 731 // CHECK: error: instruction not supported on this GPU
 732
 733 tbuffer_load_format_d16_x v0, off, s[0:3]
 734 // CHECK: error: instruction not supported on this GPU
 735
 736 tbuffer_load_format_d16_xy v0, off, s[0:3], format:22, 0
 737 // CHECK: error: instruction not supported on this GPU
 738
 739 tbuffer_load_format_d16_xyz v[1:2], off, s[4:7], dfmt:15, nfmt:2, s1
 740 // CHECK: error: instruction not supported on this GPU
 741
 742 tbuffer_load_format_d16_xyzw v[0:1], off, s[0:3], format:22, 0
 743 // CHECK: error: instruction not supported on this GPU
 744
 745 tbuffer_store_format_d16_x v0, v1, s[4:7], format:33, 0 idxen
 746 // CHECK: error: instruction not supported on this GPU
 747
 748 tbuffer_store_format_d16_xy v0, v1, s[4:7], format:33, 0 idxen
 749 // CHECK: error: instruction not supported on this GPU
 750
 751 tbuffer_store_format_d16_xyz v[1:2], off, s[4:7], dfmt:15, nfmt:2, s1
 752 // CHECK: error: instruction not supported on this GPU
 753
 754 tbuffer_store_format_d16_xyzw v[0:1], v2, s[4:7], format:33, 0 idxen
 755 // CHECK: error: instruction not supported on this GPU
 756
 757 v_accvgpr_read_b32 a0, a0
 758 // CHECK: error: instruction not supported on this GPU
 759
 760 v_accvgpr_write_b32 a0, 65
 761 // CHECK: error: instruction not supported on this GPU
 762
 763 v_add3_u32 v1, v2, v3, v4
 764 // CHECK: error: instruction not supported on this GPU
 765
 766 v_add_co_ci_u32 v1, sext(v1), sext(v4) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:DWORD
 767 // CHECK: error: instruction not supported on this GPU
 768
 769 v_add_co_ci_u32_dpp v0, vcc, v0, v0, vcc dpp8:[7,6,5,4,3,2,1,0] fi:1
 770 // CHECK: error: instruction not supported on this GPU
 771
 772 v_add_co_ci_u32_e32 v255, vcc, v1, v2, vcc
 773 // CHECK: error: instruction not supported on this GPU
 774
 775 v_add_co_ci_u32_e64 v255, s12, v1, v2, s6
 776 // CHECK: error: instruction not supported on this GPU
 777
 778 v_add_co_ci_u32_sdwa v1, v1, v4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:DWORD
 779 // CHECK: error: instruction not supported on this GPU
 780
 781 v_add_f16 v0, s[0:1], v0
 782 // CHECK: error: instruction not supported on this GPU
 783
 784 v_add_f16_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 785 // CHECK: error: instruction not supported on this GPU
 786
 787 v_add_f16_e32 v1, 64.0, v2
 788 // CHECK: error: instruction not supported on this GPU
 789
 790 v_add_f16_e64 v0, 0x3456, v0
 791 // CHECK: error: instruction not supported on this GPU
 792
 793 v_add_f16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
 794 // CHECK: error: instruction not supported on this GPU
 795
 796 v_add_i16 v255, v1, v2
 797 // CHECK: error: instruction not supported on this GPU
 798
 799 v_add_lshl_u32 v1, v2, v3, v4
 800 // CHECK: error: instruction not supported on this GPU
 801
 802 v_add_nc_i16 v255, v1, v2
 803 // CHECK: error: instruction not supported on this GPU
 804
 805 v_add_nc_i32 v255, v1, v2
 806 // CHECK: error: instruction not supported on this GPU
 807
 808 v_add_nc_u16 v255, v1, v2
 809 // CHECK: error: instruction not supported on this GPU
 810
 811 v_add_nc_u32_dpp v5, v1, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1
 812 // CHECK: error: instruction not supported on this GPU
 813
 814 v_add_nc_u32_e32 v255, v1, v2
 815 // CHECK: error: instruction not supported on this GPU
 816
 817 v_add_nc_u32_e64 v255, v1, v2
 818 // CHECK: error: instruction not supported on this GPU
 819
 820 v_add_nc_u32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
 821 // CHECK: error: instruction not supported on this GPU
 822
 823 v_add_u16 v0, (i1+100)*2, v0
 824 // CHECK: error: instruction not supported on this GPU
 825
 826 v_add_u16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 827 // CHECK: error: instruction not supported on this GPU
 828
 829 v_add_u16_e64 v255, v1, v2
 830 // CHECK: error: instruction not supported on this GPU
 831
 832 v_add_u16_sdwa v0, scc, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
 833 // CHECK: error: instruction not supported on this GPU
 834
 835 v_add_u32 v0, execz, v0
 836 // CHECK: error: instruction not supported on this GPU
 837
 838 v_add_u32_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 839 // CHECK: error: instruction not supported on this GPU
 840
 841 v_add_u32_e32 v1, s1, v3
 842 // CHECK: error: instruction not supported on this GPU
 843
 844 v_add_u32_e64 v0, scc, v0
 845 // CHECK: error: instruction not supported on this GPU
 846
 847 v_add_u32_sdwa v1, vcc, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
 848 // CHECK: error: instruction not supported on this GPU
 849
 850 v_addc_co_u32 v0, vcc, shared_base, v0, vcc
 851 // CHECK: error: instruction not supported on this GPU
 852
 853 v_addc_co_u32_dpp v255, vcc, v1, v2, vcc quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 854 // CHECK: error: instruction not supported on this GPU
 855
 856 v_addc_co_u32_e32 v3, vcc, 12345, v3, vcc
 857 // CHECK: error: instruction not supported on this GPU
 858
 859 v_addc_co_u32_e64 v255, s[12:13], v1, v2, s[6:7]
 860 // CHECK: error: instruction not supported on this GPU
 861
 862 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
 863 // CHECK: error: instruction not supported on this GPU
 864
 865 v_and_or_b32 v1, v2, v3, v4
 866 // CHECK: error: instruction not supported on this GPU
 867
 868 v_ashrrev_i16 v0, lds_direct, v0
 869 // CHECK: error: instruction not supported on this GPU
 870
 871 v_ashrrev_i16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 872 // CHECK: error: instruction not supported on this GPU
 873
 874 v_ashrrev_i16_e64 v255, v1, v2
 875 // CHECK: error: instruction not supported on this GPU
 876
 877 v_ashrrev_i16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
 878 // CHECK: error: instruction not supported on this GPU
 879
 880 v_ashrrev_i64 v[0:1], 0x100, s[0:1]
 881 // CHECK: error: instruction not supported on this GPU
 882
 883 v_ceil_f16 v0, -0.5
 884 // CHECK: error: instruction not supported on this GPU
 885
 886 v_ceil_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 887 // CHECK: error: instruction not supported on this GPU
 888
 889 v_ceil_f16_e32 v255, v1
 890 // CHECK: error: instruction not supported on this GPU
 891
 892 v_ceil_f16_e64 v0, -|v1|
 893 // CHECK: error: instruction not supported on this GPU
 894
 895 v_ceil_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
 896 // CHECK: error: instruction not supported on this GPU
 897
 898 v_cmp_class_f16 vcc, -1, v2
 899 // CHECK: error: instruction not supported on this GPU
 900
 901 v_cmp_class_f16_e64 flat_scratch, v1, v2
 902 // CHECK: error: instruction not supported on this GPU
 903
 904 v_cmp_class_f16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 905 // CHECK: error: instruction not supported on this GPU
 906
 907 v_cmp_eq_f16 vcc, -1, v0
 908 // CHECK: error: instruction not supported on this GPU
 909
 910 v_cmp_eq_f16_e64 flat_scratch, v1, v2
 911 // CHECK: error: instruction not supported on this GPU
 912
 913 v_cmp_eq_f16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 914 // CHECK: error: instruction not supported on this GPU
 915
 916 v_cmp_eq_i16 vcc, -1, v2
 917 // CHECK: error: instruction not supported on this GPU
 918
 919 v_cmp_eq_i16_e64 flat_scratch, v1, v2
 920 // CHECK: error: instruction not supported on this GPU
 921
 922 v_cmp_eq_i16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 923 // CHECK: error: instruction not supported on this GPU
 924
 925 v_cmp_eq_u16 vcc, -1, v2
 926 // CHECK: error: instruction not supported on this GPU
 927
 928 v_cmp_eq_u16_e64 flat_scratch, v1, v2
 929 // CHECK: error: instruction not supported on this GPU
 930
 931 v_cmp_eq_u16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 932 // CHECK: error: instruction not supported on this GPU
 933
 934 v_cmp_f_f16 vcc, -1, v2
 935 // CHECK: error: instruction not supported on this GPU
 936
 937 v_cmp_f_f16_e64 flat_scratch, v1, v2
 938 // CHECK: error: instruction not supported on this GPU
 939
 940 v_cmp_f_f16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 941 // CHECK: error: instruction not supported on this GPU
 942
 943 v_cmp_f_i16 vcc, -1, v2
 944 // CHECK: error: instruction not supported on this GPU
 945
 946 v_cmp_f_i16_e64 flat_scratch, v1, v2
 947 // CHECK: error: instruction not supported on this GPU
 948
 949 v_cmp_f_i16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 950 // CHECK: error: instruction not supported on this GPU
 951
 952 v_cmp_f_u16 vcc, -1, v2
 953 // CHECK: error: instruction not supported on this GPU
 954
 955 v_cmp_f_u16_e64 flat_scratch, v1, v2
 956 // CHECK: error: instruction not supported on this GPU
 957
 958 v_cmp_f_u16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 959 // CHECK: error: instruction not supported on this GPU
 960
 961 v_cmp_ge_f16 vcc, -1, v2
 962 // CHECK: error: instruction not supported on this GPU
 963
 964 v_cmp_ge_f16_e64 flat_scratch, v1, v2
 965 // CHECK: error: instruction not supported on this GPU
 966
 967 v_cmp_ge_f16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 968 // CHECK: error: instruction not supported on this GPU
 969
 970 v_cmp_ge_i16 vcc, -1, v2
 971 // CHECK: error: instruction not supported on this GPU
 972
 973 v_cmp_ge_i16_e64 flat_scratch, v1, v2
 974 // CHECK: error: instruction not supported on this GPU
 975
 976 v_cmp_ge_i16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 977 // CHECK: error: instruction not supported on this GPU
 978
 979 v_cmp_ge_u16 vcc, -1, v2
 980 // CHECK: error: instruction not supported on this GPU
 981
 982 v_cmp_ge_u16_e64 flat_scratch, v1, v2
 983 // CHECK: error: instruction not supported on this GPU
 984
 985 v_cmp_ge_u16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 986 // CHECK: error: instruction not supported on this GPU
 987
 988 v_cmp_gt_f16 vcc, -1, v2
 989 // CHECK: error: instruction not supported on this GPU
 990
 991 v_cmp_gt_f16_e64 flat_scratch, v1, v2
 992 // CHECK: error: instruction not supported on this GPU
 993
 994 v_cmp_gt_f16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 995 // CHECK: error: instruction not supported on this GPU
 996
 997 v_cmp_gt_i16 vcc, -1, v2
 998 // CHECK: error: instruction not supported on this GPU
 999
1000 v_cmp_gt_i16_e64 flat_scratch, v1, v2
1001 // CHECK: error: instruction not supported on this GPU
1002
1003 v_cmp_gt_i16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1004 // CHECK: error: instruction not supported on this GPU
1005
1006 v_cmp_gt_u16 vcc, -1, v2
1007 // CHECK: error: instruction not supported on this GPU
1008
1009 v_cmp_gt_u16_e64 flat_scratch, v1, v2
1010 // CHECK: error: instruction not supported on this GPU
1011
1012 v_cmp_gt_u16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1013 // CHECK: error: instruction not supported on this GPU
1014
1015 v_cmp_le_f16 vcc, -1, v2
1016 // CHECK: error: instruction not supported on this GPU
1017
1018 v_cmp_le_f16_e64 flat_scratch, v1, v2
1019 // CHECK: error: instruction not supported on this GPU
1020
1021 v_cmp_le_f16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1022 // CHECK: error: instruction not supported on this GPU
1023
1024 v_cmp_le_i16 vcc, -1, v2
1025 // CHECK: error: instruction not supported on this GPU
1026
1027 v_cmp_le_i16_e64 flat_scratch, v1, v2
1028 // CHECK: error: instruction not supported on this GPU
1029
1030 v_cmp_le_i16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1031 // CHECK: error: instruction not supported on this GPU
1032
1033 v_cmp_le_u16 vcc, -1, v2
1034 // CHECK: error: instruction not supported on this GPU
1035
1036 v_cmp_le_u16_e64 flat_scratch, v1, v2
1037 // CHECK: error: instruction not supported on this GPU
1038
1039 v_cmp_le_u16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1040 // CHECK: error: instruction not supported on this GPU
1041
1042 v_cmp_lg_f16 vcc, -1, v2
1043 // CHECK: error: instruction not supported on this GPU
1044
1045 v_cmp_lg_f16_e64 flat_scratch, v1, v2
1046 // CHECK: error: instruction not supported on this GPU
1047
1048 v_cmp_lg_f16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1049 // CHECK: error: instruction not supported on this GPU
1050
1051 v_cmp_lt_f16 vcc, -1, v2
1052 // CHECK: error: instruction not supported on this GPU
1053
1054 v_cmp_lt_f16_e64 flat_scratch, v1, v2
1055 // CHECK: error: instruction not supported on this GPU
1056
1057 v_cmp_lt_f16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1058 // CHECK: error: instruction not supported on this GPU
1059
1060 v_cmp_lt_i16 vcc, -1, v2
1061 // CHECK: error: instruction not supported on this GPU
1062
1063 v_cmp_lt_i16_e64 flat_scratch, v1, v2
1064 // CHECK: error: instruction not supported on this GPU
1065
1066 v_cmp_lt_i16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1067 // CHECK: error: instruction not supported on this GPU
1068
1069 v_cmp_lt_u16 vcc, -1, v2
1070 // CHECK: error: instruction not supported on this GPU
1071
1072 v_cmp_lt_u16_e64 flat_scratch, v1, v2
1073 // CHECK: error: instruction not supported on this GPU
1074
1075 v_cmp_lt_u16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1076 // CHECK: error: instruction not supported on this GPU
1077
1078 v_cmp_ne_i16 vcc, -1, v2
1079 // CHECK: error: instruction not supported on this GPU
1080
1081 v_cmp_ne_i16_e64 flat_scratch, v1, v2
1082 // CHECK: error: instruction not supported on this GPU
1083
1084 v_cmp_ne_i16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1085 // CHECK: error: instruction not supported on this GPU
1086
1087 v_cmp_ne_u16 vcc, -1, v2
1088 // CHECK: error: instruction not supported on this GPU
1089
1090 v_cmp_ne_u16_e64 flat_scratch, v1, v2
1091 // CHECK: error: instruction not supported on this GPU
1092
1093 v_cmp_ne_u16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1094 // CHECK: error: instruction not supported on this GPU
1095
1096 v_cmp_neq_f16 vcc, -1, v2
1097 // CHECK: error: instruction not supported on this GPU
1098
1099 v_cmp_neq_f16_e64 flat_scratch, v1, v2
1100 // CHECK: error: instruction not supported on this GPU
1101
1102 v_cmp_neq_f16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1103 // CHECK: error: instruction not supported on this GPU
1104
1105 v_cmp_nge_f16 vcc, -1, v2
1106 // CHECK: error: instruction not supported on this GPU
1107
1108 v_cmp_nge_f16_e64 flat_scratch, v1, v2
1109 // CHECK: error: instruction not supported on this GPU
1110
1111 v_cmp_nge_f16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1112 // CHECK: error: instruction not supported on this GPU
1113
1114 v_cmp_ngt_f16 vcc, -1, v2
1115 // CHECK: error: instruction not supported on this GPU
1116
1117 v_cmp_ngt_f16_e64 flat_scratch, v1, v2
1118 // CHECK: error: instruction not supported on this GPU
1119
1120 v_cmp_ngt_f16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1121 // CHECK: error: instruction not supported on this GPU
1122
1123 v_cmp_nle_f16 vcc, -1, v2
1124 // CHECK: error: instruction not supported on this GPU
1125
1126 v_cmp_nle_f16_e64 flat_scratch, v1, v2
1127 // CHECK: error: instruction not supported on this GPU
1128
1129 v_cmp_nle_f16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1130 // CHECK: error: instruction not supported on this GPU
1131
1132 v_cmp_nlg_f16 vcc, -1, v2
1133 // CHECK: error: instruction not supported on this GPU
1134
1135 v_cmp_nlg_f16_e64 flat_scratch, v1, v2
1136 // CHECK: error: instruction not supported on this GPU
1137
1138 v_cmp_nlg_f16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1139 // CHECK: error: instruction not supported on this GPU
1140
1141 v_cmp_nlt_f16 vcc, -1, v2
1142 // CHECK: error: instruction not supported on this GPU
1143
1144 v_cmp_nlt_f16_e64 flat_scratch, v1, v2
1145 // CHECK: error: instruction not supported on this GPU
1146
1147 v_cmp_nlt_f16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1148 // CHECK: error: instruction not supported on this GPU
1149
1150 v_cmp_o_f16 vcc, -1, v2
1151 // CHECK: error: instruction not supported on this GPU
1152
1153 v_cmp_o_f16_e64 flat_scratch, v1, v2
1154 // CHECK: error: instruction not supported on this GPU
1155
1156 v_cmp_o_f16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1157 // CHECK: error: instruction not supported on this GPU
1158
1159 v_cmp_t_i16 vcc, -1, v2
1160 // CHECK: error: instruction not supported on this GPU
1161
1162 v_cmp_t_i16_e64 flat_scratch, v1, v2
1163 // CHECK: error: instruction not supported on this GPU
1164
1165 v_cmp_t_i16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1166 // CHECK: error: instruction not supported on this GPU
1167
1168 v_cmp_t_u16 vcc, -1, v2
1169 // CHECK: error: instruction not supported on this GPU
1170
1171 v_cmp_t_u16_e64 flat_scratch, v1, v2
1172 // CHECK: error: instruction not supported on this GPU
1173
1174 v_cmp_t_u16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1175 // CHECK: error: instruction not supported on this GPU
1176
1177 v_cmp_tru_f16 vcc, -1, v2
1178 // CHECK: error: instruction not supported on this GPU
1179
1180 v_cmp_tru_f16_e64 flat_scratch, v1, v2
1181 // CHECK: error: instruction not supported on this GPU
1182
1183 v_cmp_tru_f16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1184 // CHECK: error: instruction not supported on this GPU
1185
1186 v_cmp_u_f16 vcc, -1, v2
1187 // CHECK: error: instruction not supported on this GPU
1188
1189 v_cmp_u_f16_e64 flat_scratch, v1, v2
1190 // CHECK: error: instruction not supported on this GPU
1191
1192 v_cmp_u_f16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1193 // CHECK: error: instruction not supported on this GPU
1194
1195 v_cmpx_class_f16 -1, v2
1196 // CHECK: error: instruction not supported on this GPU
1197
1198 v_cmpx_class_f16_e64 -1, v2
1199 // CHECK: error: instruction not supported on this GPU
1200
1201 v_cmpx_class_f16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1202 // CHECK: error: instruction not supported on this GPU
1203
1204 v_cmpx_eq_f16 -1, v2
1205 // CHECK: error: instruction not supported on this GPU
1206
1207 v_cmpx_eq_f16_e64 -1, v2
1208 // CHECK: error: instruction not supported on this GPU
1209
1210 v_cmpx_eq_f16_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
1211 // CHECK: error: instruction not supported on this GPU
1212
1213 v_cmpx_eq_i16 -1, v2
1214 // CHECK: error: instruction not supported on this GPU
1215
1216 v_cmpx_eq_i16_e64 -1, v2
1217 // CHECK: error: instruction not supported on this GPU
1218
1219 v_cmpx_eq_i16_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
1220 // CHECK: error: instruction not supported on this GPU
1221
1222 v_cmpx_eq_u16 -1, v2
1223 // CHECK: error: instruction not supported on this GPU
1224
1225 v_cmpx_eq_u16_e64 -1, v2
1226 // CHECK: error: instruction not supported on this GPU
1227
1228 v_cmpx_eq_u16_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
1229 // CHECK: error: instruction not supported on this GPU
1230
1231 v_cmpx_f_f16 -1, v2
1232 // CHECK: error: instruction not supported on this GPU
1233
1234 v_cmpx_f_f16_e64 -1, v2
1235 // CHECK: error: instruction not supported on this GPU
1236
1237 v_cmpx_f_f16_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
1238 // CHECK: error: instruction not supported on this GPU
1239
1240 v_cmpx_f_i16 vcc, -1, v2
1241 // CHECK: error: instruction not supported on this GPU
1242
1243 v_cmpx_f_i16_e64 exec, v1, v2
1244 // CHECK: error: instruction not supported on this GPU
1245
1246 v_cmpx_f_i16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1247 // CHECK: error: instruction not supported on this GPU
1248
1249 v_cmpx_f_u16 vcc, -1, v2
1250 // CHECK: error: instruction not supported on this GPU
1251
1252 v_cmpx_f_u16_e64 exec, v1, v2
1253 // CHECK: error: instruction not supported on this GPU
1254
1255 v_cmpx_f_u16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1256 // CHECK: error: instruction not supported on this GPU
1257
1258 v_cmpx_ge_f16 -1, v2
1259 // CHECK: error: instruction not supported on this GPU
1260
1261 v_cmpx_ge_f16_e64 -1, v2
1262 // CHECK: error: instruction not supported on this GPU
1263
1264 v_cmpx_ge_f16_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
1265 // CHECK: error: instruction not supported on this GPU
1266
1267 v_cmpx_ge_i16 -1, v2
1268 // CHECK: error: instruction not supported on this GPU
1269
1270 v_cmpx_ge_i16_e64 -1, v2
1271 // CHECK: error: instruction not supported on this GPU
1272
1273 v_cmpx_ge_i16_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
1274 // CHECK: error: instruction not supported on this GPU
1275
1276 v_cmpx_ge_u16 -1, v2
1277 // CHECK: error: instruction not supported on this GPU
1278
1279 v_cmpx_ge_u16_e64 -1, v2
1280 // CHECK: error: instruction not supported on this GPU
1281
1282 v_cmpx_ge_u16_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
1283 // CHECK: error: instruction not supported on this GPU
1284
1285 v_cmpx_gt_f16 -1, v2
1286 // CHECK: error: instruction not supported on this GPU
1287
1288 v_cmpx_gt_f16_e64 -1, v2
1289 // CHECK: error: instruction not supported on this GPU
1290
1291 v_cmpx_gt_f16_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
1292 // CHECK: error: instruction not supported on this GPU
1293
1294 v_cmpx_gt_i16 -1, v2
1295 // CHECK: error: instruction not supported on this GPU
1296
1297 v_cmpx_gt_i16_e64 -1, v2
1298 // CHECK: error: instruction not supported on this GPU
1299
1300 v_cmpx_gt_i16_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
1301 // CHECK: error: instruction not supported on this GPU
1302
1303 v_cmpx_gt_u16 -1, v2
1304 // CHECK: error: instruction not supported on this GPU
1305
1306 v_cmpx_gt_u16_e64 -1, v2
1307 // CHECK: error: instruction not supported on this GPU
1308
1309 v_cmpx_gt_u16_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
1310 // CHECK: error: instruction not supported on this GPU
1311
1312 v_cmpx_le_f16 -1, v2
1313 // CHECK: error: instruction not supported on this GPU
1314
1315 v_cmpx_le_f16_e64 -1, v2
1316 // CHECK: error: instruction not supported on this GPU
1317
1318 v_cmpx_le_f16_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
1319 // CHECK: error: instruction not supported on this GPU
1320
1321 v_cmpx_le_i16 -1, v2
1322 // CHECK: error: instruction not supported on this GPU
1323
1324 v_cmpx_le_i16_e64 -1, v2
1325 // CHECK: error: instruction not supported on this GPU
1326
1327 v_cmpx_le_i16_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
1328 // CHECK: error: instruction not supported on this GPU
1329
1330 v_cmpx_le_u16 -1, v2
1331 // CHECK: error: instruction not supported on this GPU
1332
1333 v_cmpx_le_u16_e64 -1, v2
1334 // CHECK: error: instruction not supported on this GPU
1335
1336 v_cmpx_le_u16_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
1337 // CHECK: error: instruction not supported on this GPU
1338
1339 v_cmpx_lg_f16 -1, v2
1340 // CHECK: error: instruction not supported on this GPU
1341
1342 v_cmpx_lg_f16_e64 -1, v2
1343 // CHECK: error: instruction not supported on this GPU
1344
1345 v_cmpx_lg_f16_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
1346 // CHECK: error: instruction not supported on this GPU
1347
1348 v_cmpx_lt_f16 -1, v2
1349 // CHECK: error: instruction not supported on this GPU
1350
1351 v_cmpx_lt_f16_e64 -1, v2
1352 // CHECK: error: instruction not supported on this GPU
1353
1354 v_cmpx_lt_f16_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
1355 // CHECK: error: instruction not supported on this GPU
1356
1357 v_cmpx_lt_i16 -1, v2
1358 // CHECK: error: instruction not supported on this GPU
1359
1360 v_cmpx_lt_i16_e64 -1, v2
1361 // CHECK: error: instruction not supported on this GPU
1362
1363 v_cmpx_lt_i16_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
1364 // CHECK: error: instruction not supported on this GPU
1365
1366 v_cmpx_lt_u16 -1, v2
1367 // CHECK: error: instruction not supported on this GPU
1368
1369 v_cmpx_lt_u16_e64 -1, v2
1370 // CHECK: error: instruction not supported on this GPU
1371
1372 v_cmpx_lt_u16_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
1373 // CHECK: error: instruction not supported on this GPU
1374
1375 v_cmpx_ne_i16 -1, v2
1376 // CHECK: error: instruction not supported on this GPU
1377
1378 v_cmpx_ne_i16_e64 -1, v2
1379 // CHECK: error: instruction not supported on this GPU
1380
1381 v_cmpx_ne_i16_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
1382 // CHECK: error: instruction not supported on this GPU
1383
1384 v_cmpx_ne_u16 -1, v2
1385 // CHECK: error: instruction not supported on this GPU
1386
1387 v_cmpx_ne_u16_e64 -1, v2
1388 // CHECK: error: instruction not supported on this GPU
1389
1390 v_cmpx_ne_u16_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
1391 // CHECK: error: instruction not supported on this GPU
1392
1393 v_cmpx_neq_f16 -1, v2
1394 // CHECK: error: instruction not supported on this GPU
1395
1396 v_cmpx_neq_f16_e64 -1, v2
1397 // CHECK: error: instruction not supported on this GPU
1398
1399 v_cmpx_neq_f16_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
1400 // CHECK: error: instruction not supported on this GPU
1401
1402 v_cmpx_nge_f16 -1, v2
1403 // CHECK: error: instruction not supported on this GPU
1404
1405 v_cmpx_nge_f16_e64 -1, v2
1406 // CHECK: error: instruction not supported on this GPU
1407
1408 v_cmpx_nge_f16_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
1409 // CHECK: error: instruction not supported on this GPU
1410
1411 v_cmpx_ngt_f16 -1, v2
1412 // CHECK: error: instruction not supported on this GPU
1413
1414 v_cmpx_ngt_f16_e64 -1, v2
1415 // CHECK: error: instruction not supported on this GPU
1416
1417 v_cmpx_ngt_f16_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
1418 // CHECK: error: instruction not supported on this GPU
1419
1420 v_cmpx_nle_f16 -1, v2
1421 // CHECK: error: instruction not supported on this GPU
1422
1423 v_cmpx_nle_f16_e64 -1, v2
1424 // CHECK: error: instruction not supported on this GPU
1425
1426 v_cmpx_nle_f16_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
1427 // CHECK: error: instruction not supported on this GPU
1428
1429 v_cmpx_nlg_f16 -1, v2
1430 // CHECK: error: instruction not supported on this GPU
1431
1432 v_cmpx_nlg_f16_e64 -1, v2
1433 // CHECK: error: instruction not supported on this GPU
1434
1435 v_cmpx_nlg_f16_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
1436 // CHECK: error: instruction not supported on this GPU
1437
1438 v_cmpx_nlt_f16 -1, v2
1439 // CHECK: error: instruction not supported on this GPU
1440
1441 v_cmpx_nlt_f16_e64 -1, v2
1442 // CHECK: error: instruction not supported on this GPU
1443
1444 v_cmpx_nlt_f16_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
1445 // CHECK: error: instruction not supported on this GPU
1446
1447 v_cmpx_o_f16 -1, v2
1448 // CHECK: error: instruction not supported on this GPU
1449
1450 v_cmpx_o_f16_e64 -1, v2
1451 // CHECK: error: instruction not supported on this GPU
1452
1453 v_cmpx_o_f16_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
1454 // CHECK: error: instruction not supported on this GPU
1455
1456 v_cmpx_t_i16 vcc, -1, v2
1457 // CHECK: error: instruction not supported on this GPU
1458
1459 v_cmpx_t_i16_e64 exec, v1, v2
1460 // CHECK: error: instruction not supported on this GPU
1461
1462 v_cmpx_t_i16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1463 // CHECK: error: instruction not supported on this GPU
1464
1465 v_cmpx_t_u16 vcc, -1, v2
1466 // CHECK: error: instruction not supported on this GPU
1467
1468 v_cmpx_t_u16_e64 exec, v1, v2
1469 // CHECK: error: instruction not supported on this GPU
1470
1471 v_cmpx_t_u16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
1472 // CHECK: error: instruction not supported on this GPU
1473
1474 v_cmpx_tru_f16 -1, v2
1475 // CHECK: error: instruction not supported on this GPU
1476
1477 v_cmpx_tru_f16_e64 -1, v2
1478 // CHECK: error: instruction not supported on this GPU
1479
1480 v_cmpx_tru_f16_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
1481 // CHECK: error: instruction not supported on this GPU
1482
1483 v_cmpx_u_f16 -1, v2
1484 // CHECK: error: instruction not supported on this GPU
1485
1486 v_cmpx_u_f16_e64 -1, v2
1487 // CHECK: error: instruction not supported on this GPU
1488
1489 v_cmpx_u_f16_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
1490 // CHECK: error: instruction not supported on this GPU
1491
1492 v_cos_f16 v1, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
1493 // CHECK: error: instruction not supported on this GPU
1494
1495 v_cos_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1496 // CHECK: error: instruction not supported on this GPU
1497
1498 v_cos_f16_e32 v255, v1
1499 // CHECK: error: instruction not supported on this GPU
1500
1501 v_cos_f16_e64 v255, v1
1502 // CHECK: error: instruction not supported on this GPU
1503
1504 v_cos_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1505 // CHECK: error: instruction not supported on this GPU
1506
1507 v_cvt_f16_i16 v1, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
1508 // CHECK: error: instruction not supported on this GPU
1509
1510 v_cvt_f16_i16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1511 // CHECK: error: instruction not supported on this GPU
1512
1513 v_cvt_f16_i16_e32 v255, v1
1514 // CHECK: error: instruction not supported on this GPU
1515
1516 v_cvt_f16_i16_e64 v255, v1
1517 // CHECK: error: instruction not supported on this GPU
1518
1519 v_cvt_f16_i16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1520 // CHECK: error: instruction not supported on this GPU
1521
1522 v_cvt_f16_u16 v0, src_lds_direct
1523 // CHECK: error: instruction not supported on this GPU
1524
1525 v_cvt_f16_u16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1526 // CHECK: error: instruction not supported on this GPU
1527
1528 v_cvt_f16_u16_e32 v255, v1
1529 // CHECK: error: instruction not supported on this GPU
1530
1531 v_cvt_f16_u16_e64 v255, v1
1532 // CHECK: error: instruction not supported on this GPU
1533
1534 v_cvt_f16_u16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1535 // CHECK: error: instruction not supported on this GPU
1536
1537 v_cvt_i16_f16 v1, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
1538 // CHECK: error: instruction not supported on this GPU
1539
1540 v_cvt_i16_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1541 // CHECK: error: instruction not supported on this GPU
1542
1543 v_cvt_i16_f16_e32 v255, v1
1544 // CHECK: error: instruction not supported on this GPU
1545
1546 v_cvt_i16_f16_e64 v255, v1
1547 // CHECK: error: instruction not supported on this GPU
1548
1549 v_cvt_i16_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1550 // CHECK: error: instruction not supported on this GPU
1551
1552 v_cvt_norm_i16_f16 v5, -4.0
1553 // CHECK: error: instruction not supported on this GPU
1554
1555 v_cvt_norm_i16_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1556 // CHECK: error: instruction not supported on this GPU
1557
1558 v_cvt_norm_i16_f16_e32 v255, v1
1559 // CHECK: error: instruction not supported on this GPU
1560
1561 v_cvt_norm_i16_f16_e64 v255, v1
1562 // CHECK: error: instruction not supported on this GPU
1563
1564 v_cvt_norm_i16_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1565 // CHECK: error: instruction not supported on this GPU
1566
1567 v_cvt_norm_u16_f16 v5, s101
1568 // CHECK: error: instruction not supported on this GPU
1569
1570 v_cvt_norm_u16_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1571 // CHECK: error: instruction not supported on this GPU
1572
1573 v_cvt_norm_u16_f16_e32 v255, v1
1574 // CHECK: error: instruction not supported on this GPU
1575
1576 v_cvt_norm_u16_f16_e64 v255, v1
1577 // CHECK: error: instruction not supported on this GPU
1578
1579 v_cvt_norm_u16_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1580 // CHECK: error: instruction not supported on this GPU
1581
1582 v_cvt_pknorm_i16_f16 v255, v1, v2
1583 // CHECK: error: instruction not supported on this GPU
1584
1585 v_cvt_pknorm_u16_f16 v255, v1, v2
1586 // CHECK: error: instruction not supported on this GPU
1587
1588 v_cvt_u16_f16 v1, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
1589 // CHECK: error: instruction not supported on this GPU
1590
1591 v_cvt_u16_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1592 // CHECK: error: instruction not supported on this GPU
1593
1594 v_cvt_u16_f16_e32 v255, v1
1595 // CHECK: error: instruction not supported on this GPU
1596
1597 v_cvt_u16_f16_e64 v255, v1
1598 // CHECK: error: instruction not supported on this GPU
1599
1600 v_cvt_u16_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1601 // CHECK: error: instruction not supported on this GPU
1602
1603 v_div_fixup_f16 v255, v1, v2, v3
1604 // CHECK: error: instruction not supported on this GPU
1605
1606 v_div_fixup_f16_e64 v5, 0.5, v2, v3
1607 // CHECK: error: instruction not supported on this GPU
1608
1609 v_div_fixup_legacy_f16 v255, v1, v2, v3
1610 // CHECK: error: instruction not supported on this GPU
1611
1612 v_div_fixup_legacy_f16_e64 v5, 0.5, v2, v3
1613 // CHECK: error: instruction not supported on this GPU
1614
1615 v_dot2_f32_f16 v0, -v1, -v2, -v3
1616 // CHECK: error: instruction not supported on this GPU
1617
1618 v_dot2_i32_i16 v0, -v1, -v2, -v3
1619 // CHECK: error: instruction not supported on this GPU
1620
1621 v_dot2_u32_u16 v0, -v1, -v2, -v3
1622 // CHECK: error: instruction not supported on this GPU
1623
1624 v_dot2c_f32_f16 v0, v1, v2
1625 // CHECK: error: instruction not supported on this GPU
1626
1627 v_dot2c_f32_f16_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1628 // CHECK: error: instruction not supported on this GPU
1629
1630 v_dot2c_f32_f16_e32 v255, v1, v2
1631 // CHECK: error: instruction not supported on this GPU
1632
1633 v_dot2c_i32_i16 v0, v1, v2
1634 // CHECK: error: instruction not supported on this GPU
1635
1636 v_dot2c_i32_i16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1637 // CHECK: error: instruction not supported on this GPU
1638
1639 v_dot4_i32_i8 v0, v1, v2, v3
1640 // CHECK: error: instruction not supported on this GPU
1641
1642 v_dot4_u32_u8 v0, v1, v2, v3
1643 // CHECK: error: instruction not supported on this GPU
1644
1645 v_dot4c_i32_i8 v0, v1, v2
1646 // CHECK: error: instruction not supported on this GPU
1647
1648 v_dot4c_i32_i8_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1649 // CHECK: error: instruction not supported on this GPU
1650
1651 v_dot4c_i32_i8_e32 v255, v1, v2
1652 // CHECK: error: instruction not supported on this GPU
1653
1654 v_dot8_i32_i4 v0, v1, v2, v3
1655 // CHECK: error: instruction not supported on this GPU
1656
1657 v_dot8_u32_u4 v0, v1, v2, v3
1658 // CHECK: error: instruction not supported on this GPU
1659
1660 v_dot8c_i32_i4 v0, v1, v2
1661 // CHECK: error: instruction not supported on this GPU
1662
1663 v_dot8c_i32_i4_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1664 // CHECK: error: instruction not supported on this GPU
1665
1666 v_exp_f16 v1, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
1667 // CHECK: error: instruction not supported on this GPU
1668
1669 v_exp_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1670 // CHECK: error: instruction not supported on this GPU
1671
1672 v_exp_f16_e32 v255, v1
1673 // CHECK: error: instruction not supported on this GPU
1674
1675 v_exp_f16_e64 v255, v1
1676 // CHECK: error: instruction not supported on this GPU
1677
1678 v_exp_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1679 // CHECK: error: instruction not supported on this GPU
1680
1681 v_floor_f16 v1, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
1682 // CHECK: error: instruction not supported on this GPU
1683
1684 v_floor_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1685 // CHECK: error: instruction not supported on this GPU
1686
1687 v_floor_f16_e32 v255, v1
1688 // CHECK: error: instruction not supported on this GPU
1689
1690 v_floor_f16_e64 v255, v1
1691 // CHECK: error: instruction not supported on this GPU
1692
1693 v_floor_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1694 // CHECK: error: instruction not supported on this GPU
1695
1696 v_fma_f16 v255, v1, v2, v3
1697 // CHECK: error: instruction not supported on this GPU
1698
1699 v_fma_f16_e64 v5, v1, v2, v3
1700 // CHECK: error: instruction not supported on this GPU
1701
1702 v_fma_legacy_f16 v255, v1, v2, v3
1703 // CHECK: error: instruction not supported on this GPU
1704
1705 v_fma_legacy_f16_e64 v5, v1, v2, v3
1706 // CHECK: error: instruction not supported on this GPU
1707
1708 v_fma_mix_f32 v0, -abs(v1), v2, v3
1709 // CHECK: error: instruction not supported on this GPU
1710
1711 v_fma_mixhi_f16 v0, -v1, abs(v2), -abs(v3)
1712 // CHECK: error: instruction not supported on this GPU
1713
1714 v_fma_mixlo_f16 v0, abs(v1), -v2, abs(v3)
1715 // CHECK: error: instruction not supported on this GPU
1716
1717 v_fmaak_f32 v255, v1, v2, 0x1121
1718 // CHECK: error: instruction not supported on this GPU
1719
1720 v_fmac_f16 v5, 0x1234, v2
1721 // CHECK: error: instruction not supported on this GPU
1722
1723 v_fmac_f16_dpp v5, v1, v2 quad_perm:[3,2,1,0] row_mask:0x0 bank_mask:0x0
1724 // CHECK: error: instruction not supported on this GPU
1725
1726 v_fmac_f16_e32 v255, v1, v2
1727 // CHECK: error: instruction not supported on this GPU
1728
1729 v_fmac_f16_e64 v255, v1, v2
1730 // CHECK: error: instruction not supported on this GPU
1731
1732 v_fmac_f32 v0, v1, v2
1733 // CHECK: error: instruction not supported on this GPU
1734
1735 v_fmac_f32_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1736 // CHECK: error: instruction not supported on this GPU
1737
1738 v_fmac_f32_e32 v255, v1, v2
1739 // CHECK: error: instruction not supported on this GPU
1740
1741 v_fmac_f32_e64 v255, v1, v2
1742 // CHECK: error: instruction not supported on this GPU
1743
1744 v_fmamk_f32 v255, v1, 0x1121, v3
1745 // CHECK: error: instruction not supported on this GPU
1746
1747 v_fract_f16 v1, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
1748 // CHECK: error: instruction not supported on this GPU
1749
1750 v_fract_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1751 // CHECK: error: instruction not supported on this GPU
1752
1753 v_fract_f16_e32 v255, v1
1754 // CHECK: error: instruction not supported on this GPU
1755
1756 v_fract_f16_e64 v255, v1
1757 // CHECK: error: instruction not supported on this GPU
1758
1759 v_fract_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1760 // CHECK: error: instruction not supported on this GPU
1761
1762 v_frexp_exp_i16_f16 v1, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
1763 // CHECK: error: instruction not supported on this GPU
1764
1765 v_frexp_exp_i16_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1766 // CHECK: error: instruction not supported on this GPU
1767
1768 v_frexp_exp_i16_f16_e32 v255, v1
1769 // CHECK: error: instruction not supported on this GPU
1770
1771 v_frexp_exp_i16_f16_e64 v255, v1
1772 // CHECK: error: instruction not supported on this GPU
1773
1774 v_frexp_exp_i16_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1775 // CHECK: error: instruction not supported on this GPU
1776
1777 v_frexp_mant_f16 v1, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
1778 // CHECK: error: instruction not supported on this GPU
1779
1780 v_frexp_mant_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1781 // CHECK: error: instruction not supported on this GPU
1782
1783 v_frexp_mant_f16_e32 v255, v1
1784 // CHECK: error: instruction not supported on this GPU
1785
1786 v_frexp_mant_f16_e64 v255, v1
1787 // CHECK: error: instruction not supported on this GPU
1788
1789 v_frexp_mant_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1790 // CHECK: error: instruction not supported on this GPU
1791
1792 v_interp_p1ll_f16 v255, v2, attr0.x
1793 // CHECK: error: instruction not supported on this GPU
1794
1795 v_interp_p1lv_f16 v255, v2, attr0.x, v3
1796 // CHECK: error: instruction not supported on this GPU
1797
1798 v_interp_p2_f16 v255, v2, attr0.x, v3
1799 // CHECK: error: instruction not supported on this GPU
1800
1801 v_interp_p2_legacy_f16 v255, v2, attr0.x, v3
1802 // CHECK: error: instruction not supported on this GPU
1803
1804 v_ldexp_f16 v1, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
1805 // CHECK: error: instruction not supported on this GPU
1806
1807 v_ldexp_f16_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1808 // CHECK: error: instruction not supported on this GPU
1809
1810 v_ldexp_f16_e32 v255, v1, v2
1811 // CHECK: error: instruction not supported on this GPU
1812
1813 v_ldexp_f16_e64 v255, v1, v2
1814 // CHECK: error: instruction not supported on this GPU
1815
1816 v_ldexp_f16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1817 // CHECK: error: instruction not supported on this GPU
1818
1819 v_log_f16 v1, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
1820 // CHECK: error: instruction not supported on this GPU
1821
1822 v_log_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1823 // CHECK: error: instruction not supported on this GPU
1824
1825 v_log_f16_e32 v255, v1
1826 // CHECK: error: instruction not supported on this GPU
1827
1828 v_log_f16_e64 v255, v1
1829 // CHECK: error: instruction not supported on this GPU
1830
1831 v_log_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
1832 // CHECK: error: instruction not supported on this GPU
1833
1834 v_lshl_add_u32 v1, v2, v3, v4
1835 // CHECK: error: instruction not supported on this GPU
1836
1837 v_lshl_or_b32 v1, v2, v3, v4
1838 // CHECK: error: instruction not supported on this GPU
1839
1840 v_lshlrev_b16 v0, lds_direct, v0
1841 // CHECK: error: instruction not supported on this GPU
1842
1843 v_lshlrev_b16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1844 // CHECK: error: instruction not supported on this GPU
1845
1846 v_lshlrev_b16_e64 v255, v1, v2
1847 // CHECK: error: instruction not supported on this GPU
1848
1849 v_lshlrev_b16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1850 // CHECK: error: instruction not supported on this GPU
1851
1852 v_lshlrev_b64 v[254:255], v1, v[2:3]
1853 // CHECK: error: instruction not supported on this GPU
1854
1855 v_lshrrev_b16 v0, lds_direct, v0
1856 // CHECK: error: instruction not supported on this GPU
1857
1858 v_lshrrev_b16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1859 // CHECK: error: instruction not supported on this GPU
1860
1861 v_lshrrev_b16_e64 v255, v1, v2
1862 // CHECK: error: instruction not supported on this GPU
1863
1864 v_lshrrev_b16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1865 // CHECK: error: instruction not supported on this GPU
1866
1867 v_lshrrev_b64 v[254:255], v1, v[2:3]
1868 // CHECK: error: instruction not supported on this GPU
1869
1870 v_mac_f16 v1, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
1871 // CHECK: error: instruction not supported on this GPU
1872
1873 v_mac_f16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1874 // CHECK: error: instruction not supported on this GPU
1875
1876 v_mac_f16_e64 v0, -4.0, flat_scratch_lo
1877 // CHECK: error: instruction not supported on this GPU
1878
1879 v_mac_f16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1880 // CHECK: error: instruction not supported on this GPU
1881
1882 v_mad_f16 v255, v1, v2, v3
1883 // CHECK: error: instruction not supported on this GPU
1884
1885 v_mad_f16_e64 v5, 0.5, v2, v3
1886 // CHECK: error: instruction not supported on this GPU
1887
1888 v_mad_i16 v255, v1, v2, v3
1889 // CHECK: error: instruction not supported on this GPU
1890
1891 v_mad_i16_e64 v5, -1, v2, v3
1892 // CHECK: error: instruction not supported on this GPU
1893
1894 v_mad_i32_i16 v255, v1, v2, v3
1895 // CHECK: error: instruction not supported on this GPU
1896
1897 v_mad_legacy_f16 v255, v1, v2, v3
1898 // CHECK: error: instruction not supported on this GPU
1899
1900 v_mad_legacy_f16_e64 v5, 0.5, v2, v3
1901 // CHECK: error: instruction not supported on this GPU
1902
1903 v_mad_legacy_i16 v255, v1, v2, v3
1904 // CHECK: error: instruction not supported on this GPU
1905
1906 v_mad_legacy_i16_e64 v5, 0, v2, v3
1907 // CHECK: error: instruction not supported on this GPU
1908
1909 v_mad_legacy_u16 v255, v1, v2, v3
1910 // CHECK: error: instruction not supported on this GPU
1911
1912 v_mad_legacy_u16_e64 v5, 0, v2, v3
1913 // CHECK: error: instruction not supported on this GPU
1914
1915 v_mad_mix_f32 v0, -abs(v1), v2, v3
1916 // CHECK: error: instruction not supported on this GPU
1917
1918 v_mad_mixhi_f16 v0, -v1, abs(v2), -abs(v3)
1919 // CHECK: error: instruction not supported on this GPU
1920
1921 v_mad_mixlo_f16 v0, abs(v1), -v2, abs(v3)
1922 // CHECK: error: instruction not supported on this GPU
1923
1924 v_mad_u16 v255, v1, v2, v3
1925 // CHECK: error: instruction not supported on this GPU
1926
1927 v_mad_u16_e64 v5, -1, v2, v3
1928 // CHECK: error: instruction not supported on this GPU
1929
1930 v_mad_u32_u16 v255, v1, v2, v3
1931 // CHECK: error: instruction not supported on this GPU
1932
1933 v_madak_f16 v0, src_lds_direct, v0, 0x1121
1934 // CHECK: error: instruction not supported on this GPU
1935
1936 v_madmk_f16 v0, src_lds_direct, 0x1121, v0
1937 // CHECK: error: instruction not supported on this GPU
1938
1939 v_max3_f16 v0, src_lds_direct, v0, v0
1940 // CHECK: error: instruction not supported on this GPU
1941
1942 v_max3_i16 v1, v2, v3, v4
1943 // CHECK: error: instruction not supported on this GPU
1944
1945 v_max3_u16 v1, v2, v3, v4
1946 // CHECK: error: instruction not supported on this GPU
1947
1948 v_max_f16 v0, execz, v0
1949 // CHECK: error: instruction not supported on this GPU
1950
1951 v_max_f16_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1952 // CHECK: error: instruction not supported on this GPU
1953
1954 v_max_f16_e32 v255, v1, v2
1955 // CHECK: error: instruction not supported on this GPU
1956
1957 v_max_f16_e64 v255, v1, v2
1958 // CHECK: error: instruction not supported on this GPU
1959
1960 v_max_f16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1961 // CHECK: error: instruction not supported on this GPU
1962
1963 v_max_i16 v1, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
1964 // CHECK: error: instruction not supported on this GPU
1965
1966 v_max_i16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1967 // CHECK: error: instruction not supported on this GPU
1968
1969 v_max_i16_e64 v255, v1, v2
1970 // CHECK: error: instruction not supported on this GPU
1971
1972 v_max_i16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1973 // CHECK: error: instruction not supported on this GPU
1974
1975 v_max_u16 v1, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
1976 // CHECK: error: instruction not supported on this GPU
1977
1978 v_max_u16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1979 // CHECK: error: instruction not supported on this GPU
1980
1981 v_max_u16_e64 v255, v1, v2
1982 // CHECK: error: instruction not supported on this GPU
1983
1984 v_max_u16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1985 // CHECK: error: instruction not supported on this GPU
1986
1987 v_med3_f16 v1, v2, v3, v4
1988 // CHECK: error: instruction not supported on this GPU
1989
1990 v_med3_i16 v1, v2, v3, v4
1991 // CHECK: error: instruction not supported on this GPU
1992
1993 v_med3_u16 v1, v2, v3, v4
1994 // CHECK: error: instruction not supported on this GPU
1995
1996 v_mfma_f32_16x16x16f16 a[0:3], a[0:1], a[1:2], -2.0
1997 // CHECK: error: instruction not supported on this GPU
1998
1999 v_mfma_f32_16x16x1f32 a[0:15], a0, a1, -2.0
2000 // CHECK: error: instruction not supported on this GPU
2001
2002 v_mfma_f32_16x16x2bf16 a[0:15], a0, a1, -2.0
2003 // CHECK: error: instruction not supported on this GPU
2004
2005 v_mfma_f32_16x16x4f16 a[0:15], a[0:1], a[1:2], -2.0
2006 // CHECK: error: instruction not supported on this GPU
2007
2008 v_mfma_f32_16x16x4f32 a[0:3], a0, a1, -2.0
2009 // CHECK: error: instruction not supported on this GPU
2010
2011 v_mfma_f32_16x16x8bf16 a[0:3], a0, a1, -2.0
2012 // CHECK: error: instruction not supported on this GPU
2013
2014 v_mfma_f32_32x32x1f32 a[0:31], 1, v1, a[1:32]
2015 // CHECK: error: instruction not supported on this GPU
2016
2017 v_mfma_f32_32x32x2bf16 a[0:31], a0, a1, -2.0
2018 // CHECK: error: instruction not supported on this GPU
2019
2020 v_mfma_f32_32x32x2f32 a[0:15], a0, a1, -2.0
2021 // CHECK: error: instruction not supported on this GPU
2022
2023 v_mfma_f32_32x32x4bf16 a[0:15], a0, a1, -2.0
2024 // CHECK: error: instruction not supported on this GPU
2025
2026 v_mfma_f32_32x32x4f16 a[0:31], a[0:1], a[1:2], -2.0
2027 // CHECK: error: instruction not supported on this GPU
2028
2029 v_mfma_f32_32x32x8f16 a[0:15], a[0:1], a[1:2], -2.0
2030 // CHECK: error: instruction not supported on this GPU
2031
2032 v_mfma_f32_4x4x1f32 a[0:3], a0, a1, -2.0
2033 // CHECK: error: instruction not supported on this GPU
2034
2035 v_mfma_f32_4x4x2bf16 a[0:3], a0, a1, -2.0
2036 // CHECK: error: instruction not supported on this GPU
2037
2038 v_mfma_f32_4x4x4f16 a[0:3], a[0:1], a[1:2], -2.0
2039 // CHECK: error: instruction not supported on this GPU
2040
2041 v_mfma_i32_16x16x16i8 a[0:3], a0, a1, 2
2042 // CHECK: error: instruction not supported on this GPU
2043
2044 v_mfma_i32_16x16x4i8 a[0:15], a0, a1, 2
2045 // CHECK: error: instruction not supported on this GPU
2046
2047 v_mfma_i32_32x32x4i8 a[0:31], a0, a1, 2
2048 // CHECK: error: instruction not supported on this GPU
2049
2050 v_mfma_i32_32x32x8i8 a[0:15], a0, a1, 2
2051 // CHECK: error: instruction not supported on this GPU
2052
2053 v_mfma_i32_4x4x4i8 a[0:3], a0, a1, 2
2054 // CHECK: error: instruction not supported on this GPU
2055
2056 v_min3_f16 v1, v2, v3, v4
2057 // CHECK: error: instruction not supported on this GPU
2058
2059 v_min3_i16 v0, src_lds_direct, v0, v0
2060 // CHECK: error: instruction not supported on this GPU
2061
2062 v_min3_u16 v1, v2, v3, v4
2063 // CHECK: error: instruction not supported on this GPU
2064
2065 v_min_f16 v1, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
2066 // CHECK: error: instruction not supported on this GPU
2067
2068 v_min_f16_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2069 // CHECK: error: instruction not supported on this GPU
2070
2071 v_min_f16_e32 v255, v1, v2
2072 // CHECK: error: instruction not supported on this GPU
2073
2074 v_min_f16_e64 v255, v1, v2
2075 // CHECK: error: instruction not supported on this GPU
2076
2077 v_min_f16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
2078 // CHECK: error: instruction not supported on this GPU
2079
2080 v_min_i16 v1, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
2081 // CHECK: error: instruction not supported on this GPU
2082
2083 v_min_i16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2084 // CHECK: error: instruction not supported on this GPU
2085
2086 v_min_i16_e64 v255, v1, v2
2087 // CHECK: error: instruction not supported on this GPU
2088
2089 v_min_i16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
2090 // CHECK: error: instruction not supported on this GPU
2091
2092 v_min_u16 v1, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
2093 // CHECK: error: instruction not supported on this GPU
2094
2095 v_min_u16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2096 // CHECK: error: instruction not supported on this GPU
2097
2098 v_min_u16_e64 v255, v1, v2
2099 // CHECK: error: instruction not supported on this GPU
2100
2101 v_min_u16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
2102 // CHECK: error: instruction not supported on this GPU
2103
2104 v_movrelsd_2_b32 v0, v255 dpp8:[7,6,5,4,3,2,1,0]
2105 // CHECK: error: instruction not supported on this GPU
2106
2107 v_movrelsd_2_b32_dpp v0, v2 quad_perm:[3,2,1,0] row_mask:0x0 bank_mask:0x0
2108 // CHECK: error: instruction not supported on this GPU
2109
2110 v_movrelsd_2_b32_e32 v5, 1
2111 // CHECK: error: instruction not supported on this GPU
2112
2113 v_movrelsd_2_b32_e64 v255, v1
2114 // CHECK: error: instruction not supported on this GPU
2115
2116 v_movrelsd_2_b32_sdwa v0, 0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2117 // CHECK: error: instruction not supported on this GPU
2118
2119 v_mul_f16 v1, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
2120 // CHECK: error: instruction not supported on this GPU
2121
2122 v_mul_f16_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2123 // CHECK: error: instruction not supported on this GPU
2124
2125 v_mul_f16_e32 v255, v1, v2
2126 // CHECK: error: instruction not supported on this GPU
2127
2128 v_mul_f16_e64 v255, v1, v2
2129 // CHECK: error: instruction not supported on this GPU
2130
2131 v_mul_f16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
2132 // CHECK: error: instruction not supported on this GPU
2133
2134 v_mul_lo_u16 v1, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
2135 // CHECK: error: instruction not supported on this GPU
2136
2137 v_mul_lo_u16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2138 // CHECK: error: instruction not supported on this GPU
2139
2140 v_mul_lo_u16_e64 v255, v1, v2
2141 // CHECK: error: instruction not supported on this GPU
2142
2143 v_mul_lo_u16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
2144 // CHECK: error: instruction not supported on this GPU
2145
2146 v_or3_b32 v1, v2, v3, v4
2147 // CHECK: error: instruction not supported on this GPU
2148
2149 v_pack_b32_f16 v1, v2, v3
2150 // CHECK: error: instruction not supported on this GPU
2151
2152 v_perm_b32 v255, v1, v2, v3
2153 // CHECK: error: instruction not supported on this GPU
2154
2155 v_permlane16_b32 v0, lds_direct, s0, s0
2156 // CHECK: error: instruction not supported on this GPU
2157
2158 v_permlanex16_b32 v0, lds_direct, s0, s0
2159 // CHECK: error: instruction not supported on this GPU
2160
2161 v_pipeflush
2162 // CHECK: error: instruction not supported on this GPU
2163
2164 v_pipeflush_e64
2165 // CHECK: error: instruction not supported on this GPU
2166
2167 v_pk_add_f16 v0, execz, v0
2168 // CHECK: error: instruction not supported on this GPU
2169
2170 v_pk_add_i16 v0, src_lds_direct, v0
2171 // CHECK: error: instruction not supported on this GPU
2172
2173 v_pk_add_u16 v0, v1, v2
2174 // CHECK: error: instruction not supported on this GPU
2175
2176 v_pk_ashrrev_i16 v0, lds_direct, v0
2177 // CHECK: error: instruction not supported on this GPU
2178
2179 v_pk_fma_f16 v0, v1, v2, v3
2180 // CHECK: error: instruction not supported on this GPU
2181
2182 v_pk_fmac_f16 v0, v1, v2
2183 // CHECK: error: instruction not supported on this GPU
2184
2185 v_pk_lshlrev_b16 v0, lds_direct, v0
2186 // CHECK: error: instruction not supported on this GPU
2187
2188 v_pk_lshrrev_b16 v0, lds_direct, v0
2189 // CHECK: error: instruction not supported on this GPU
2190
2191 v_pk_mad_i16 v0, src_lds_direct, v0, v0
2192 // CHECK: error: instruction not supported on this GPU
2193
2194 v_pk_mad_u16 v255, v1, v2, v3
2195 // CHECK: error: instruction not supported on this GPU
2196
2197 v_pk_max_f16 v0, v1, v2
2198 // CHECK: error: instruction not supported on this GPU
2199
2200 v_pk_max_i16 v0, v1, v2
2201 // CHECK: error: instruction not supported on this GPU
2202
2203 v_pk_max_u16 v0, v1, v2
2204 // CHECK: error: instruction not supported on this GPU
2205
2206 v_pk_min_f16 v0, v1, v2
2207 // CHECK: error: instruction not supported on this GPU
2208
2209 v_pk_min_i16 v0, v1, v2
2210 // CHECK: error: instruction not supported on this GPU
2211
2212 v_pk_min_u16 v0, v1, v2
2213 // CHECK: error: instruction not supported on this GPU
2214
2215 v_pk_mul_f16 v0, v1, v2
2216 // CHECK: error: instruction not supported on this GPU
2217
2218 v_pk_mul_lo_u16 v0, v1, v2
2219 // CHECK: error: instruction not supported on this GPU
2220
2221 v_pk_sub_i16 v0, v1, v2
2222 // CHECK: error: instruction not supported on this GPU
2223
2224 v_pk_sub_u16 v255, v1, v2
2225 // CHECK: error: instruction not supported on this GPU
2226
2227 v_rcp_f16 v1, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
2228 // CHECK: error: instruction not supported on this GPU
2229
2230 v_rcp_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2231 // CHECK: error: instruction not supported on this GPU
2232
2233 v_rcp_f16_e32 v255, v1
2234 // CHECK: error: instruction not supported on this GPU
2235
2236 v_rcp_f16_e64 v255, v1
2237 // CHECK: error: instruction not supported on this GPU
2238
2239 v_rcp_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2240 // CHECK: error: instruction not supported on this GPU
2241
2242 v_rndne_f16 v1, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
2243 // CHECK: error: instruction not supported on this GPU
2244
2245 v_rndne_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2246 // CHECK: error: instruction not supported on this GPU
2247
2248 v_rndne_f16_e32 v255, v1
2249 // CHECK: error: instruction not supported on this GPU
2250
2251 v_rndne_f16_e64 v255, v1
2252 // CHECK: error: instruction not supported on this GPU
2253
2254 v_rndne_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2255 // CHECK: error: instruction not supported on this GPU
2256
2257 v_rsq_f16 v1, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
2258 // CHECK: error: instruction not supported on this GPU
2259
2260 v_rsq_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2261 // CHECK: error: instruction not supported on this GPU
2262
2263 v_rsq_f16_e32 v255, v1
2264 // CHECK: error: instruction not supported on this GPU
2265
2266 v_rsq_f16_e64 v255, v1
2267 // CHECK: error: instruction not supported on this GPU
2268
2269 v_rsq_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2270 // CHECK: error: instruction not supported on this GPU
2271
2272 v_sat_pk_u8_i16 v255, v1
2273 // CHECK: error: instruction not supported on this GPU
2274
2275 v_sat_pk_u8_i16_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0x0 bank_mask:0x0
2276 // CHECK: error: instruction not supported on this GPU
2277
2278 v_sat_pk_u8_i16_e64 v5, -1
2279 // CHECK: error: instruction not supported on this GPU
2280
2281 v_sat_pk_u8_i16_sdwa v5, sext(v1) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2282 // CHECK: error: instruction not supported on this GPU
2283
2284 v_screen_partition_4se_b32 v5, -1
2285 // CHECK: error: instruction not supported on this GPU
2286
2287 v_screen_partition_4se_b32_dpp v5, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0 bound_ctrl:0
2288 // CHECK: error: instruction not supported on this GPU
2289
2290 v_screen_partition_4se_b32_e64 v5, -1
2291 // CHECK: error: instruction not supported on this GPU
2292
2293 v_screen_partition_4se_b32_sdwa v5, v1 src0_sel:BYTE_0
2294 // CHECK: error: instruction not supported on this GPU
2295
2296 v_sin_f16 v1, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
2297 // CHECK: error: instruction not supported on this GPU
2298
2299 v_sin_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2300 // CHECK: error: instruction not supported on this GPU
2301
2302 v_sin_f16_e32 v255, v1
2303 // CHECK: error: instruction not supported on this GPU
2304
2305 v_sin_f16_e64 v255, v1
2306 // CHECK: error: instruction not supported on this GPU
2307
2308 v_sin_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2309 // CHECK: error: instruction not supported on this GPU
2310
2311 v_sqrt_f16 v1, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
2312 // CHECK: error: instruction not supported on this GPU
2313
2314 v_sqrt_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2315 // CHECK: error: instruction not supported on this GPU
2316
2317 v_sqrt_f16_e32 v255, v1
2318 // CHECK: error: instruction not supported on this GPU
2319
2320 v_sqrt_f16_e64 v255, v1
2321 // CHECK: error: instruction not supported on this GPU
2322
2323 v_sqrt_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2324 // CHECK: error: instruction not supported on this GPU
2325
2326 v_sub_co_ci_u32_dpp v0, vcc, v0, v0, vcc dpp8:[7,6,5,4,3,2,1,0] fi:1
2327 // CHECK: error: instruction not supported on this GPU
2328
2329 v_sub_co_ci_u32_e32 v255, vcc, v1, v2, vcc
2330 // CHECK: error: instruction not supported on this GPU
2331
2332 v_sub_co_ci_u32_e64 v255, s12, v1, v2, s6
2333 // CHECK: error: instruction not supported on this GPU
2334
2335 v_sub_co_ci_u32_sdwa v1, v1, v4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:DWORD
2336 // CHECK: error: instruction not supported on this GPU
2337
2338 v_sub_f16 v1, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
2339 // CHECK: error: instruction not supported on this GPU
2340
2341 v_sub_f16_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2342 // CHECK: error: instruction not supported on this GPU
2343
2344 v_sub_f16_e32 v255, v1, v2
2345 // CHECK: error: instruction not supported on this GPU
2346
2347 v_sub_f16_e64 v255, v1, v2
2348 // CHECK: error: instruction not supported on this GPU
2349
2350 v_sub_f16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
2351 // CHECK: error: instruction not supported on this GPU
2352
2353 v_sub_i16 v255, v1, v2
2354 // CHECK: error: instruction not supported on this GPU
2355
2356 v_sub_nc_i16 v255, v1, v2
2357 // CHECK: error: instruction not supported on this GPU
2358
2359 v_sub_nc_i32 v255, v1, v2
2360 // CHECK: error: instruction not supported on this GPU
2361
2362 v_sub_nc_u16 v255, v1, v2
2363 // CHECK: error: instruction not supported on this GPU
2364
2365 v_sub_nc_u32_dpp v5, v1, v2 dpp8:[7,6,5,4,3,2,1,0]
2366 // CHECK: error: instruction not supported on this GPU
2367
2368 v_sub_nc_u32_e32 v255, v1, v2
2369 // CHECK: error: instruction not supported on this GPU
2370
2371 v_sub_nc_u32_e64 v255, v1, v2
2372 // CHECK: error: instruction not supported on this GPU
2373
2374 v_sub_nc_u32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
2375 // CHECK: error: instruction not supported on this GPU
2376
2377 v_sub_u16 v1, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
2378 // CHECK: error: instruction not supported on this GPU
2379
2380 v_sub_u16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2381 // CHECK: error: instruction not supported on this GPU
2382
2383 v_sub_u16_e64 v255, v1, v2
2384 // CHECK: error: instruction not supported on this GPU
2385
2386 v_sub_u16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
2387 // CHECK: error: instruction not supported on this GPU
2388
2389 v_sub_u32 v1, 4.0, v2
2390 // CHECK: error: instruction not supported on this GPU
2391
2392 v_sub_u32_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2393 // CHECK: error: instruction not supported on this GPU
2394
2395 v_sub_u32_e32 v1, s1, v3
2396 // CHECK: error: instruction not supported on this GPU
2397
2398 v_sub_u32_e64 v255, s[12:13], v1, v2
2399 // CHECK: error: instruction not supported on this GPU
2400
2401 v_sub_u32_sdwa v1, vcc, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
2402 // CHECK: error: instruction not supported on this GPU
2403
2404 v_subb_co_u32 v1, vcc, v2, v3, vcc row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
2405 // CHECK: error: instruction not supported on this GPU
2406
2407 v_subb_co_u32_dpp v255, vcc, v1, v2, vcc quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2408 // CHECK: error: instruction not supported on this GPU
2409
2410 v_subb_co_u32_e64 v255, s[12:13], v1, v2, s[6:7]
2411 // CHECK: error: instruction not supported on this GPU
2412
2413 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
2414 // CHECK: error: instruction not supported on this GPU
2415
2416 v_subbrev_co_u32 v0, vcc, src_lds_direct, v0, vcc
2417 // CHECK: error: instruction not supported on this GPU
2418
2419 v_subbrev_co_u32_dpp v255, vcc, v1, v2, vcc quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2420 // CHECK: error: instruction not supported on this GPU
2421
2422 v_subbrev_co_u32_e64 v255, s[12:13], v1, v2, s[6:7]
2423 // CHECK: error: instruction not supported on this GPU
2424
2425 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
2426 // CHECK: error: instruction not supported on this GPU
2427
2428 v_subrev_co_ci_u32 v0, vcc_lo, src_lds_direct, v0, vcc_lo
2429 // CHECK: error: instruction not supported on this GPU
2430
2431 v_subrev_co_ci_u32_dpp v0, vcc, v0, v0, vcc dpp8:[7,6,5,4,3,2,1,0]
2432 // CHECK: error: instruction not supported on this GPU
2433
2434 v_subrev_co_ci_u32_e32 v1, 0, v1
2435 // CHECK: error: instruction not supported on this GPU
2436
2437 v_subrev_co_ci_u32_e64 v255, s12, v1, v2, s6
2438 // CHECK: error: instruction not supported on this GPU
2439
2440 v_subrev_co_ci_u32_sdwa v1, v1, v4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:DWORD
2441 // CHECK: error: instruction not supported on this GPU
2442
2443 v_subrev_f16 v0, src_lds_direct, v0
2444 // CHECK: error: instruction not supported on this GPU
2445
2446 v_subrev_f16_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2447 // CHECK: error: instruction not supported on this GPU
2448
2449 v_subrev_f16_e32 v255, v1, v2
2450 // CHECK: error: instruction not supported on this GPU
2451
2452 v_subrev_f16_e64 v255, v1, v2
2453 // CHECK: error: instruction not supported on this GPU
2454
2455 v_subrev_f16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
2456 // CHECK: error: instruction not supported on this GPU
2457
2458 v_subrev_nc_u32 v0, src_lds_direct, v0
2459 // CHECK: error: instruction not supported on this GPU
2460
2461 v_subrev_nc_u32_dpp v5, v1, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1
2462 // CHECK: error: instruction not supported on this GPU
2463
2464 v_subrev_nc_u32_e32 v255, v1, v2
2465 // CHECK: error: instruction not supported on this GPU
2466
2467 v_subrev_nc_u32_e64 v255, v1, v2
2468 // CHECK: error: instruction not supported on this GPU
2469
2470 v_subrev_nc_u32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
2471 // CHECK: error: instruction not supported on this GPU
2472
2473 v_subrev_u16 v0, src_lds_direct, v0
2474 // CHECK: error: instruction not supported on this GPU
2475
2476 v_subrev_u16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2477 // CHECK: error: instruction not supported on this GPU
2478
2479 v_subrev_u16_e64 v255, v1, v2
2480 // CHECK: error: instruction not supported on this GPU
2481
2482 v_subrev_u16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
2483 // CHECK: error: instruction not supported on this GPU
2484
2485 v_subrev_u32 v0, src_lds_direct, v0
2486 // CHECK: error: instruction not supported on this GPU
2487
2488 v_subrev_u32_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2489 // CHECK: error: instruction not supported on this GPU
2490
2491 v_subrev_u32_e32 v1, s1, v3
2492 // CHECK: error: instruction not supported on this GPU
2493
2494 v_subrev_u32_e64 v255, s[12:13], v1, v2
2495 // CHECK: error: instruction not supported on this GPU
2496
2497 v_subrev_u32_sdwa v1, vcc, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
2498 // CHECK: error: instruction not supported on this GPU
2499
2500 v_swap_b32 v1, 1
2501 // CHECK: error: instruction not supported on this GPU
2502
2503 v_swap_b32_e32 v1, v2
2504 // CHECK: error: instruction not supported on this GPU
2505
2506 v_swaprel_b32 v255, v1
2507 // CHECK: error: instruction not supported on this GPU
2508
2509 v_trunc_f16 v1, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
2510 // CHECK: error: instruction not supported on this GPU
2511
2512 v_trunc_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2513 // CHECK: error: instruction not supported on this GPU
2514
2515 v_trunc_f16_e32 v255, v1
2516 // CHECK: error: instruction not supported on this GPU
2517
2518 v_trunc_f16_e64 v255, v1
2519 // CHECK: error: instruction not supported on this GPU
2520
2521 v_trunc_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2522 // CHECK: error: instruction not supported on this GPU
2523
2524 v_xad_u32 v1, v2, v3, v4
2525 // CHECK: error: instruction not supported on this GPU
2526
2527 v_xnor_b32 v0, v1, v2
2528 // CHECK: error: instruction not supported on this GPU
2529
2530 v_xnor_b32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2531 // CHECK: error: instruction not supported on this GPU
2532
2533 v_xnor_b32_e32 v255, v1, v2
2534 // CHECK: error: instruction not supported on this GPU
2535
2536 v_xnor_b32_e64 v255, v1, v2
2537 // CHECK: error: instruction not supported on this GPU
2538
2539 v_xnor_b32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
2540 // CHECK: error: instruction not supported on this GPU
2541
2542 v_xor3_b32 v255, v1, v2, v3
2543 // CHECK: error: instruction not supported on this GPU
2544
2545 //===----------------------------------------------------------------------===//
2546 // Unsupported e64 variants.
2547 //===----------------------------------------------------------------------===//
2548
2549 v_interp_mov_f32_e64 v255, p10, attr0.x
2550 // CHECK: error: e64 variant of this instruction is not supported
2551
2552 v_interp_p1_f32_e64 v255, v2, attr0.x
2553 // CHECK: error: e64 variant of this instruction is not supported
2554
2555 v_interp_p2_f32_e64 v255, v2, attr0.x
2556 // CHECK: error: e64 variant of this instruction is not supported
2557
2558 //===----------------------------------------------------------------------===//
2559 // Unsupported dpp variants.
2560 //===----------------------------------------------------------------------===//
2561
2562 v_add_f32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2563 // CHECK: error: dpp variant of this instruction is not supported
2564
2565 v_addc_u32_dpp v255, vcc, v1, v2, vcc quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2566 // CHECK: error: dpp variant of this instruction is not supported
2567
2568 v_and_b32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2569 // CHECK: error: dpp variant of this instruction is not supported
2570
2571 v_ashrrev_i32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2572 // CHECK: error: dpp variant of this instruction is not supported
2573
2574 v_bfrev_b32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2575 // CHECK: error: dpp variant of this instruction is not supported
2576
2577 v_ceil_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2578 // CHECK: error: dpp variant of this instruction is not supported
2579
2580 v_cos_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2581 // CHECK: error: dpp variant of this instruction is not supported
2582
2583 v_cvt_f16_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2584 // CHECK: error: dpp variant of this instruction is not supported
2585
2586 v_cvt_f32_f16_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2587 // CHECK: error: dpp variant of this instruction is not supported
2588
2589 v_cvt_f32_i32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2590 // CHECK: error: dpp variant of this instruction is not supported
2591
2592 v_cvt_f32_u32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2593 // CHECK: error: dpp variant of this instruction is not supported
2594
2595 v_cvt_f32_ubyte0_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2596 // CHECK: error: dpp variant of this instruction is not supported
2597
2598 v_cvt_f32_ubyte1_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2599 // CHECK: error: dpp variant of this instruction is not supported
2600
2601 v_cvt_f32_ubyte2_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2602 // CHECK: error: dpp variant of this instruction is not supported
2603
2604 v_cvt_f32_ubyte3_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2605 // CHECK: error: dpp variant of this instruction is not supported
2606
2607 v_cvt_flr_i32_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2608 // CHECK: error: dpp variant of this instruction is not supported
2609
2610 v_cvt_i32_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2611 // CHECK: error: dpp variant of this instruction is not supported
2612
2613 v_cvt_off_f32_i4_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2614 // CHECK: error: dpp variant of this instruction is not supported
2615
2616 v_cvt_rpi_i32_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2617 // CHECK: error: dpp variant of this instruction is not supported
2618
2619 v_cvt_u32_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2620 // CHECK: error: dpp variant of this instruction is not supported
2621
2622 v_exp_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2623 // CHECK: error: dpp variant of this instruction is not supported
2624
2625 v_exp_legacy_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2626 // CHECK: error: dpp variant of this instruction is not supported
2627
2628 v_ffbh_i32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2629 // CHECK: error: dpp variant of this instruction is not supported
2630
2631 v_ffbh_u32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2632 // CHECK: error: dpp variant of this instruction is not supported
2633
2634 v_ffbl_b32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2635 // CHECK: error: dpp variant of this instruction is not supported
2636
2637 v_floor_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2638 // CHECK: error: dpp variant of this instruction is not supported
2639
2640 v_fract_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2641 // CHECK: error: dpp variant of this instruction is not supported
2642
2643 v_frexp_exp_i32_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2644 // CHECK: error: dpp variant of this instruction is not supported
2645
2646 v_frexp_mant_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2647 // CHECK: error: dpp variant of this instruction is not supported
2648
2649 v_log_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2650 // CHECK: error: dpp variant of this instruction is not supported
2651
2652 v_log_legacy_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2653 // CHECK: error: dpp variant of this instruction is not supported
2654
2655 v_lshlrev_b32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2656 // CHECK: error: dpp variant of this instruction is not supported
2657
2658 v_lshrrev_b32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2659 // CHECK: error: dpp variant of this instruction is not supported
2660
2661 v_mac_f32_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2662 // CHECK: error: dpp variant of this instruction is not supported
2663
2664 v_max_f32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2665 // CHECK: error: dpp variant of this instruction is not supported
2666
2667 v_max_i32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2668 // CHECK: error: dpp variant of this instruction is not supported
2669
2670 v_max_u32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2671 // CHECK: error: dpp variant of this instruction is not supported
2672
2673 v_min_f32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2674 // CHECK: error: dpp variant of this instruction is not supported
2675
2676 v_min_i32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2677 // CHECK: error: dpp variant of this instruction is not supported
2678
2679 v_min_u32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2680 // CHECK: error: dpp variant of this instruction is not supported
2681
2682 v_mov_b32_dpp v0, v1 row_bcast:15 row_mask:0x1 bank_mask:0x1
2683 // CHECK: error: dpp variant of this instruction is not supported
2684
2685 v_movreld_b32_dpp v1, v0 quad_perm:[3,2,1,0] row_mask:0x0 bank_mask:0x0
2686 // CHECK: error: dpp variant of this instruction is not supported
2687
2688 v_movrels_b32_dpp v1, v0 quad_perm:[3,2,1,0] row_mask:0x0 bank_mask:0x0 fi:1
2689 // CHECK: error: dpp variant of this instruction is not supported
2690
2691 v_movrelsd_b32_dpp v0, v255 quad_perm:[3,2,1,0] row_mask:0x0 bank_mask:0x0
2692 // CHECK: error: dpp variant of this instruction is not supported
2693
2694 v_mul_f32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2695 // CHECK: error: dpp variant of this instruction is not supported
2696
2697 v_mul_hi_i32_i24_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2698 // CHECK: error: dpp variant of this instruction is not supported
2699
2700 v_mul_hi_u32_u24_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2701 // CHECK: error: dpp variant of this instruction is not supported
2702
2703 v_mul_i32_i24_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2704 // CHECK: error: dpp variant of this instruction is not supported
2705
2706 v_mul_legacy_f32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2707 // CHECK: error: dpp variant of this instruction is not supported
2708
2709 v_mul_u32_u24_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2710 // CHECK: error: dpp variant of this instruction is not supported
2711
2712 v_not_b32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2713 // CHECK: error: dpp variant of this instruction is not supported
2714
2715 v_or_b32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2716 // CHECK: error: dpp variant of this instruction is not supported
2717
2718 v_rcp_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2719 // CHECK: error: dpp variant of this instruction is not supported
2720
2721 v_rcp_iflag_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2722 // CHECK: error: dpp variant of this instruction is not supported
2723
2724 v_rndne_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2725 // CHECK: error: dpp variant of this instruction is not supported
2726
2727 v_rsq_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2728 // CHECK: error: dpp variant of this instruction is not supported
2729
2730 v_sin_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2731 // CHECK: error: dpp variant of this instruction is not supported
2732
2733 v_sqrt_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2734 // CHECK: error: dpp variant of this instruction is not supported
2735
2736 v_sub_f32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2737 // CHECK: error: dpp variant of this instruction is not supported
2738
2739 v_subb_u32_dpp v255, vcc, v1, v2, vcc quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2740 // CHECK: error: dpp variant of this instruction is not supported
2741
2742 v_subbrev_u32_dpp v255, vcc, v1, v2, vcc quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2743 // CHECK: error: dpp variant of this instruction is not supported
2744
2745 v_subrev_f32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2746 // CHECK: error: dpp variant of this instruction is not supported
2747
2748 v_trunc_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2749 // CHECK: error: dpp variant of this instruction is not supported
2750
2751 v_xor_b32_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
2752 // CHECK: error: dpp variant of this instruction is not supported
2753
2754 //===----------------------------------------------------------------------===//
2755 // Unsupported sdwa variants.
2756 //===----------------------------------------------------------------------===//
2757
2758 v_add_f32_sdwa v0, v0, v0 dst_unused:UNUSED_PAD src0_sel:WORD_1
2759 // CHECK: error: sdwa variant of this instruction is not supported
2760
2761 v_addc_u32_sdwa v1, vcc, v2, v3, vcc dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
2762 // CHECK: error: sdwa variant of this instruction is not supported
2763
2764 v_and_b32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
2765 // CHECK: error: sdwa variant of this instruction is not supported
2766
2767 v_ashrrev_i32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
2768 // CHECK: error: sdwa variant of this instruction is not supported
2769
2770 v_bfrev_b32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2771 // CHECK: error: sdwa variant of this instruction is not supported
2772
2773 v_ceil_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2774 // CHECK: error: sdwa variant of this instruction is not supported
2775
2776 v_cmp_class_f32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2777 // CHECK: error: sdwa variant of this instruction is not supported
2778
2779 v_cmp_eq_f32_sdwa exec, s2, v2 src0_sel:WORD_1 src1_sel:BYTE_2
2780 // CHECK: error: sdwa variant of this instruction is not supported
2781
2782 v_cmp_eq_i32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2783 // CHECK: error: sdwa variant of this instruction is not supported
2784
2785 v_cmp_eq_u32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2786 // CHECK: error: sdwa variant of this instruction is not supported
2787
2788 v_cmp_f_f32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2789 // CHECK: error: sdwa variant of this instruction is not supported
2790
2791 v_cmp_f_i32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2792 // CHECK: error: sdwa variant of this instruction is not supported
2793
2794 v_cmp_f_u32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2795 // CHECK: error: sdwa variant of this instruction is not supported
2796
2797 v_cmp_ge_f32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2798 // CHECK: error: sdwa variant of this instruction is not supported
2799
2800 v_cmp_ge_i32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2801 // CHECK: error: sdwa variant of this instruction is not supported
2802
2803 v_cmp_ge_u32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2804 // CHECK: error: sdwa variant of this instruction is not supported
2805
2806 v_cmp_gt_f32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2807 // CHECK: error: sdwa variant of this instruction is not supported
2808
2809 v_cmp_gt_i32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2810 // CHECK: error: sdwa variant of this instruction is not supported
2811
2812 v_cmp_gt_u32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2813 // CHECK: error: sdwa variant of this instruction is not supported
2814
2815 v_cmp_le_f32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2816 // CHECK: error: sdwa variant of this instruction is not supported
2817
2818 v_cmp_le_i32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2819 // CHECK: error: sdwa variant of this instruction is not supported
2820
2821 v_cmp_le_u32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2822 // CHECK: error: sdwa variant of this instruction is not supported
2823
2824 v_cmp_lg_f32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2825 // CHECK: error: sdwa variant of this instruction is not supported
2826
2827 v_cmp_lt_f32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2828 // CHECK: error: sdwa variant of this instruction is not supported
2829
2830 v_cmp_lt_i32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2831 // CHECK: error: sdwa variant of this instruction is not supported
2832
2833 v_cmp_lt_u32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2834 // CHECK: error: sdwa variant of this instruction is not supported
2835
2836 v_cmp_ne_i32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2837 // CHECK: error: sdwa variant of this instruction is not supported
2838
2839 v_cmp_ne_u32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2840 // CHECK: error: sdwa variant of this instruction is not supported
2841
2842 v_cmp_neq_f32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2843 // CHECK: error: sdwa variant of this instruction is not supported
2844
2845 v_cmp_nge_f32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2846 // CHECK: error: sdwa variant of this instruction is not supported
2847
2848 v_cmp_ngt_f32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2849 // CHECK: error: sdwa variant of this instruction is not supported
2850
2851 v_cmp_nle_f32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2852 // CHECK: error: sdwa variant of this instruction is not supported
2853
2854 v_cmp_nlg_f32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2855 // CHECK: error: sdwa variant of this instruction is not supported
2856
2857 v_cmp_nlt_f32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2858 // CHECK: error: sdwa variant of this instruction is not supported
2859
2860 v_cmp_o_f32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2861 // CHECK: error: sdwa variant of this instruction is not supported
2862
2863 v_cmp_t_i32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2864 // CHECK: error: sdwa variant of this instruction is not supported
2865
2866 v_cmp_t_u32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2867 // CHECK: error: sdwa variant of this instruction is not supported
2868
2869 v_cmp_tru_f32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2870 // CHECK: error: sdwa variant of this instruction is not supported
2871
2872 v_cmp_u_f32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2873 // CHECK: error: sdwa variant of this instruction is not supported
2874
2875 v_cmpx_class_f32_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
2876 // CHECK: error: sdwa variant of this instruction is not supported
2877
2878 v_cmpx_eq_f32_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
2879 // CHECK: error: sdwa variant of this instruction is not supported
2880
2881 v_cmpx_eq_i32_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
2882 // CHECK: error: sdwa variant of this instruction is not supported
2883
2884 v_cmpx_eq_u32_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
2885 // CHECK: error: sdwa variant of this instruction is not supported
2886
2887 v_cmpx_f_f32_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
2888 // CHECK: error: sdwa variant of this instruction is not supported
2889
2890 v_cmpx_f_i32_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
2891 // CHECK: error: sdwa variant of this instruction is not supported
2892
2893 v_cmpx_f_u32_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
2894 // CHECK: error: sdwa variant of this instruction is not supported
2895
2896 v_cmpx_ge_f32_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
2897 // CHECK: error: sdwa variant of this instruction is not supported
2898
2899 v_cmpx_ge_i32_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
2900 // CHECK: error: sdwa variant of this instruction is not supported
2901
2902 v_cmpx_ge_u32_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
2903 // CHECK: error: sdwa variant of this instruction is not supported
2904
2905 v_cmpx_gt_f32_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
2906 // CHECK: error: sdwa variant of this instruction is not supported
2907
2908 v_cmpx_gt_i32_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
2909 // CHECK: error: sdwa variant of this instruction is not supported
2910
2911 v_cmpx_gt_u32_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
2912 // CHECK: error: sdwa variant of this instruction is not supported
2913
2914 v_cmpx_le_f32_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
2915 // CHECK: error: sdwa variant of this instruction is not supported
2916
2917 v_cmpx_le_i32_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
2918 // CHECK: error: sdwa variant of this instruction is not supported
2919
2920 v_cmpx_le_u32_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
2921 // CHECK: error: sdwa variant of this instruction is not supported
2922
2923 v_cmpx_lg_f32_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
2924 // CHECK: error: sdwa variant of this instruction is not supported
2925
2926 v_cmpx_lt_f32_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
2927 // CHECK: error: sdwa variant of this instruction is not supported
2928
2929 v_cmpx_lt_i32_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
2930 // CHECK: error: sdwa variant of this instruction is not supported
2931
2932 v_cmpx_lt_u32_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
2933 // CHECK: error: sdwa variant of this instruction is not supported
2934
2935 v_cmpx_ne_i32_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
2936 // CHECK: error: sdwa variant of this instruction is not supported
2937
2938 v_cmpx_ne_u32_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
2939 // CHECK: error: sdwa variant of this instruction is not supported
2940
2941 v_cmpx_neq_f32_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
2942 // CHECK: error: sdwa variant of this instruction is not supported
2943
2944 v_cmpx_nge_f32_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
2945 // CHECK: error: sdwa variant of this instruction is not supported
2946
2947 v_cmpx_ngt_f32_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
2948 // CHECK: error: sdwa variant of this instruction is not supported
2949
2950 v_cmpx_nle_f32_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
2951 // CHECK: error: sdwa variant of this instruction is not supported
2952
2953 v_cmpx_nlg_f32_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
2954 // CHECK: error: sdwa variant of this instruction is not supported
2955
2956 v_cmpx_nlt_f32_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
2957 // CHECK: error: sdwa variant of this instruction is not supported
2958
2959 v_cmpx_o_f32_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
2960 // CHECK: error: sdwa variant of this instruction is not supported
2961
2962 v_cmpx_t_i32_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
2963 // CHECK: error: sdwa variant of this instruction is not supported
2964
2965 v_cmpx_t_u32_sdwa exec_hi, v2 src0_sel:DWORD src1_sel:DWORD
2966 // CHECK: error: sdwa variant of this instruction is not supported
2967
2968 v_cmpx_tru_f32_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
2969 // CHECK: error: sdwa variant of this instruction is not supported
2970
2971 v_cmpx_u_f32_sdwa -v1, v2 src0_sel:DWORD src1_sel:DWORD
2972 // CHECK: error: sdwa variant of this instruction is not supported
2973
2974 v_cos_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2975 // CHECK: error: sdwa variant of this instruction is not supported
2976
2977 v_cvt_f16_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2978 // CHECK: error: sdwa variant of this instruction is not supported
2979
2980 v_cvt_f32_f16_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2981 // CHECK: error: sdwa variant of this instruction is not supported
2982
2983 v_cvt_f32_i32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2984 // CHECK: error: sdwa variant of this instruction is not supported
2985
2986 v_cvt_f32_u32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2987 // CHECK: error: sdwa variant of this instruction is not supported
2988
2989 v_cvt_f32_ubyte0_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2990 // CHECK: error: sdwa variant of this instruction is not supported
2991
2992 v_cvt_f32_ubyte1_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2993 // CHECK: error: sdwa variant of this instruction is not supported
2994
2995 v_cvt_f32_ubyte2_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2996 // CHECK: error: sdwa variant of this instruction is not supported
2997
2998 v_cvt_f32_ubyte3_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
2999 // CHECK: error: sdwa variant of this instruction is not supported
3000
3001 v_cvt_flr_i32_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3002 // CHECK: error: sdwa variant of this instruction is not supported
3003
3004 v_cvt_i32_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3005 // CHECK: error: sdwa variant of this instruction is not supported
3006
3007 v_cvt_off_f32_i4_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3008 // CHECK: error: sdwa variant of this instruction is not supported
3009
3010 v_cvt_rpi_i32_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3011 // CHECK: error: sdwa variant of this instruction is not supported
3012
3013 v_cvt_u32_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3014 // CHECK: error: sdwa variant of this instruction is not supported
3015
3016 v_exp_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3017 // CHECK: error: sdwa variant of this instruction is not supported
3018
3019 v_exp_legacy_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3020 // CHECK: error: sdwa variant of this instruction is not supported
3021
3022 v_ffbh_i32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3023 // CHECK: error: sdwa variant of this instruction is not supported
3024
3025 v_ffbh_u32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3026 // CHECK: error: sdwa variant of this instruction is not supported
3027
3028 v_ffbl_b32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3029 // CHECK: error: sdwa variant of this instruction is not supported
3030
3031 v_floor_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3032 // CHECK: error: sdwa variant of this instruction is not supported
3033
3034 v_fract_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3035 // CHECK: error: sdwa variant of this instruction is not supported
3036
3037 v_frexp_exp_i32_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3038 // CHECK: error: sdwa variant of this instruction is not supported
3039
3040 v_frexp_mant_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3041 // CHECK: error: sdwa variant of this instruction is not supported
3042
3043 v_log_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3044 // CHECK: error: sdwa variant of this instruction is not supported
3045
3046 v_log_legacy_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3047 // CHECK: error: sdwa variant of this instruction is not supported
3048
3049 v_lshlrev_b32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
3050 // CHECK: error: sdwa variant of this instruction is not supported
3051
3052 v_lshrrev_b32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
3053 // CHECK: error: sdwa variant of this instruction is not supported
3054
3055 v_mac_f32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
3056 // CHECK: error: sdwa variant of this instruction is not supported
3057
3058 v_max_f32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
3059 // CHECK: error: sdwa variant of this instruction is not supported
3060
3061 v_max_i32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
3062 // CHECK: error: sdwa variant of this instruction is not supported
3063
3064 v_max_u32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
3065 // CHECK: error: sdwa variant of this instruction is not supported
3066
3067 v_min_f32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
3068 // CHECK: error: sdwa variant of this instruction is not supported
3069
3070 v_min_i32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
3071 // CHECK: error: sdwa variant of this instruction is not supported
3072
3073 v_min_u32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
3074 // CHECK: error: sdwa variant of this instruction is not supported
3075
3076 v_mov_b32_sdwa v1, sext(-2+i1)
3077 // CHECK: error: sdwa variant of this instruction is not supported
3078
3079 v_movreld_b32_sdwa v0, 64 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3080 // CHECK: error: sdwa variant of this instruction is not supported
3081
3082 v_movrels_b32_sdwa v0, 1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3083 // CHECK: error: sdwa variant of this instruction is not supported
3084
3085 v_movrelsd_b32_sdwa v0, 1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3086 // CHECK: error: sdwa variant of this instruction is not supported
3087
3088 v_mul_f32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
3089 // CHECK: error: sdwa variant of this instruction is not supported
3090
3091 v_mul_hi_i32_i24_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
3092 // CHECK: error: sdwa variant of this instruction is not supported
3093
3094 v_mul_hi_u32_u24_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
3095 // CHECK: error: sdwa variant of this instruction is not supported
3096
3097 v_mul_i32_i24_sdwa v1, v2, v3 clamp
3098 // CHECK: error: sdwa variant of this instruction is not supported
3099
3100 v_mul_legacy_f32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
3101 // CHECK: error: sdwa variant of this instruction is not supported
3102
3103 v_mul_u32_u24_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
3104 // CHECK: error: sdwa variant of this instruction is not supported
3105
3106 v_nop_sdwa
3107 // CHECK: error: sdwa variant of this instruction is not supported
3108
3109 v_not_b32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3110 // CHECK: error: sdwa variant of this instruction is not supported
3111
3112 v_or_b32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
3113 // CHECK: error: sdwa variant of this instruction is not supported
3114
3115 v_rcp_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3116 // CHECK: error: sdwa variant of this instruction is not supported
3117
3118 v_rcp_iflag_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3119 // CHECK: error: sdwa variant of this instruction is not supported
3120
3121 v_rndne_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3122 // CHECK: error: sdwa variant of this instruction is not supported
3123
3124 v_rsq_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3125 // CHECK: error: sdwa variant of this instruction is not supported
3126
3127 v_sin_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3128 // CHECK: error: sdwa variant of this instruction is not supported
3129
3130 v_sqrt_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3131 // CHECK: error: sdwa variant of this instruction is not supported
3132
3133 v_sub_f32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
3134 // CHECK: error: sdwa variant of this instruction is not supported
3135
3136 v_subb_u32_sdwa v1, vcc, v2, v3, vcc dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
3137 // CHECK: error: sdwa variant of this instruction is not supported
3138
3139 v_subbrev_u32_sdwa v1, vcc, v2, v3, vcc dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
3140 // CHECK: error: sdwa variant of this instruction is not supported
3141
3142 v_subrev_f32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
3143 // CHECK: error: sdwa variant of this instruction is not supported
3144
3145 v_trunc_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
3146 // CHECK: error: sdwa variant of this instruction is not supported
3147
3148 v_xor_b32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
3149 // CHECK: error: sdwa variant of this instruction is not supported