llvm/test/MC/AMDGPU/gfx10_unsupported.s

   1 // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck --implicit-check-not=error: %s
   2 // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --implicit-check-not=error: %s
   3 // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1013 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --implicit-check-not=error: %s
   4
   5 //===----------------------------------------------------------------------===//
   6 // Unsupported instructions.
   7 //===----------------------------------------------------------------------===//
   8
   9 buffer_atomic_add_f32 v255, off, s[8:11], s3 offset:4095
  10 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  11
  12 buffer_atomic_pk_add_f16 v255, off, s[8:11], s3 offset:4095
  13 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  14
  15 buffer_store_lds_dword s[4:7], s0 lds
  16 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  17
  18 buffer_wbinvl1_vol
  19 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  20
  21 global_atomic_add_f32 v[1:2], v2, off
  22 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  23
  24 global_atomic_pk_add_f16 v[1:2], v2, off
  25 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  26
  27 s_cbranch_g_fork -1, s[4:5]
  28 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  29
  30 s_cbranch_i_fork exec, 12609
  31 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  32
  33 s_cbranch_join 1
  34 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  35
  36 s_dcache_inv_vol
  37 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  38
  39 s_dcache_wb_vol
  40 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  41
  42 s_rfe_restore_b64 -1, s2
  43 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  44
  45 s_set_gpr_idx_idx -1
  46 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  47
  48 s_set_gpr_idx_mode 0
  49 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  50
  51 s_set_gpr_idx_off
  52 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  53
  54 s_set_gpr_idx_on -1, 0x0
  55 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  56
  57 s_setvskip -1, s2
  58 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  59
  60 v_accvgpr_read_b32 a0, a0
  61 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  62
  63 v_accvgpr_write_b32 a0, 65
  64 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  65
  66 v_add_i16 v255, v1, v2
  67 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  68
  69 v_add_i32 lds_direct, v0, v0
  70 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  71
  72 v_add_i32_e32 v0, vcc, 0.5, v0
  73 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  74
  75 v_add_i32_e64 v1, s[0:1], v2, v3
  76 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  77
  78 v_add_u16 v0, (i1+100)*2, v0
  79 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  80
  81 v_add_u16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
  82 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  83
  84 v_add_u16_e64 v255, v1, v2
  85 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  86
  87 v_add_u16_sdwa v0, scc, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
  88 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  89
  90 v_add_u32 v0, execz, v0
  91 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  92
  93 v_add_u32_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
  94 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  95
  96 v_add_u32_e32 v1, s1, v3
  97 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
  98
  99 v_add_u32_e64 v0, scc, v0
 100 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 101
 102 v_add_u32_sdwa v1, vcc, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
 103 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 104
 105 v_addc_co_u32 v0, vcc, shared_base, v0, vcc
 106 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 107
 108 v_addc_co_u32_dpp v255, vcc, v1, v2, vcc quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 109 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 110
 111 v_addc_co_u32_e32 v3, vcc, 12345, v3, vcc
 112 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 113
 114 v_addc_co_u32_e64 v255, s[12:13], v1, v2, s[6:7]
 115 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 116
 117 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
 118 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 119
 120 v_addc_u32 v0, vcc, exec_hi, v0, vcc
 121 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 122
 123 v_addc_u32_dpp v255, vcc, v1, v2, vcc quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 124 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 125
 126 v_addc_u32_e32 v1, -1, v2, v3, s0
 127 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 128
 129 v_addc_u32_e64 v0, s[0:1], s0, s0, s[0:1]
 130 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 131
 132 v_addc_u32_sdwa v1, vcc, v2, v3, vcc dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
 133 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 134
 135 v_ashr_i32 v255, v1, v2
 136 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 137
 138 v_ashr_i32_e64 v255, v1, v2
 139 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 140
 141 v_ashr_i64 v[254:255], v[1:2], v2
 142 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 143
 144 v_cmp_f_i16 vcc, -1, v2
 145 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 146
 147 v_cmp_f_i16_e64 flat_scratch, v1, v2
 148 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 149
 150 v_cmp_f_i16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 151 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 152
 153 v_cmp_f_u16 vcc, -1, v2
 154 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 155
 156 v_cmp_f_u16_e64 flat_scratch, v1, v2
 157 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 158
 159 v_cmp_f_u16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 160 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 161
 162 v_cmp_t_i16 vcc, -1, v2
 163 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 164
 165 v_cmp_t_i16_e64 flat_scratch, v1, v2
 166 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 167
 168 v_cmp_t_i16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 169 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 170
 171 v_cmp_t_u16 vcc, -1, v2
 172 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 173
 174 v_cmp_t_u16_e64 flat_scratch, v1, v2
 175 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 176
 177 v_cmp_t_u16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 178 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 179
 180 v_cmps_eq_f32 vcc, -1, v2
 181 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 182
 183 v_cmps_eq_f32_e64 flat_scratch, v1, v2
 184 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 185
 186 v_cmps_eq_f64 vcc, -1, v[2:3]
 187 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 188
 189 v_cmps_eq_f64_e64 flat_scratch, v[1:2], v[2:3]
 190 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 191
 192 v_cmps_f_f32 vcc, -1, v2
 193 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 194
 195 v_cmps_f_f32_e64 flat_scratch, v1, v2
 196 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 197
 198 v_cmps_f_f64 vcc, -1, v[2:3]
 199 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 200
 201 v_cmps_f_f64_e64 flat_scratch, v[1:2], v[2:3]
 202 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 203
 204 v_cmps_ge_f32 vcc, -1, v2
 205 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 206
 207 v_cmps_ge_f32_e64 flat_scratch, v1, v2
 208 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 209
 210 v_cmps_ge_f64 vcc, -1, v[2:3]
 211 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 212
 213 v_cmps_ge_f64_e64 flat_scratch, v[1:2], v[2:3]
 214 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 215
 216 v_cmps_gt_f32 vcc, -1, v2
 217 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 218
 219 v_cmps_gt_f32_e64 flat_scratch, v1, v2
 220 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 221
 222 v_cmps_gt_f64 vcc, -1, v[2:3]
 223 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 224
 225 v_cmps_gt_f64_e64 flat_scratch, v[1:2], v[2:3]
 226 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 227
 228 v_cmps_le_f32 vcc, -1, v2
 229 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 230
 231 v_cmps_le_f32_e64 flat_scratch, v1, v2
 232 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 233
 234 v_cmps_le_f64 vcc, -1, v[2:3]
 235 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 236
 237 v_cmps_le_f64_e64 flat_scratch, v[1:2], v[2:3]
 238 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 239
 240 v_cmps_lg_f32 vcc, -1, v2
 241 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 242
 243 v_cmps_lg_f32_e64 flat_scratch, v1, v2
 244 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 245
 246 v_cmps_lg_f64 vcc, -1, v[2:3]
 247 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 248
 249 v_cmps_lg_f64_e64 flat_scratch, v[1:2], v[2:3]
 250 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 251
 252 v_cmps_lt_f32 vcc, -1, v2
 253 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 254
 255 v_cmps_lt_f32_e64 flat_scratch, v1, v2
 256 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 257
 258 v_cmps_lt_f64 vcc, -1, v[2:3]
 259 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 260
 261 v_cmps_lt_f64_e64 flat_scratch, v[1:2], v[2:3]
 262 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 263
 264 v_cmps_neq_f32 vcc, -1, v2
 265 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 266
 267 v_cmps_neq_f32_e64 flat_scratch, v1, v2
 268 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 269
 270 v_cmps_neq_f64 vcc, -1, v[2:3]
 271 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 272
 273 v_cmps_neq_f64_e64 flat_scratch, v[1:2], v[2:3]
 274 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 275
 276 v_cmps_nge_f32 vcc, -1, v2
 277 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 278
 279 v_cmps_nge_f32_e64 flat_scratch, v1, v2
 280 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 281
 282 v_cmps_nge_f64 vcc, -1, v[2:3]
 283 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 284
 285 v_cmps_nge_f64_e64 flat_scratch, v[1:2], v[2:3]
 286 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 287
 288 v_cmps_ngt_f32 vcc, -1, v2
 289 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 290
 291 v_cmps_ngt_f32_e64 flat_scratch, v1, v2
 292 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 293
 294 v_cmps_ngt_f64 vcc, -1, v[2:3]
 295 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 296
 297 v_cmps_ngt_f64_e64 flat_scratch, v[1:2], v[2:3]
 298 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 299
 300 v_cmps_nle_f32 vcc, -1, v2
 301 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 302
 303 v_cmps_nle_f32_e64 flat_scratch, v1, v2
 304 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 305
 306 v_cmps_nle_f64 vcc, -1, v[2:3]
 307 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 308
 309 v_cmps_nle_f64_e64 flat_scratch, v[1:2], v[2:3]
 310 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 311
 312 v_cmps_nlg_f32 vcc, -1, v2
 313 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 314
 315 v_cmps_nlg_f32_e64 flat_scratch, v1, v2
 316 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 317
 318 v_cmps_nlg_f64 vcc, -1, v[2:3]
 319 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 320
 321 v_cmps_nlg_f64_e64 flat_scratch, v[1:2], v[2:3]
 322 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 323
 324 v_cmps_nlt_f32 vcc, -1, v2
 325 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 326
 327 v_cmps_nlt_f32_e64 flat_scratch, v1, v2
 328 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 329
 330 v_cmps_nlt_f64 vcc, -1, v[2:3]
 331 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 332
 333 v_cmps_nlt_f64_e64 flat_scratch, v[1:2], v[2:3]
 334 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 335
 336 v_cmps_o_f32 vcc, -1, v2
 337 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 338
 339 v_cmps_o_f32_e64 flat_scratch, v1, v2
 340 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 341
 342 v_cmps_o_f64 vcc, -1, v[2:3]
 343 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 344
 345 v_cmps_o_f64_e64 flat_scratch, v[1:2], v[2:3]
 346 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 347
 348 v_cmps_tru_f32 vcc, -1, v2
 349 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 350
 351 v_cmps_tru_f32_e64 flat_scratch, v1, v2
 352 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 353
 354 v_cmps_tru_f64 vcc, -1, v[2:3]
 355 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 356
 357 v_cmps_tru_f64_e64 flat_scratch, v[1:2], v[2:3]
 358 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 359
 360 v_cmps_u_f32 vcc, -1, v2
 361 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 362
 363 v_cmps_u_f32_e64 flat_scratch, v1, v2
 364 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 365
 366 v_cmps_u_f64 vcc, -1, v[2:3]
 367 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 368
 369 v_cmps_u_f64_e64 flat_scratch, v[1:2], v[2:3]
 370 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 371
 372 v_cmpsx_eq_f32 vcc, -1, v2
 373 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 374
 375 v_cmpsx_eq_f32_e64 flat_scratch, v1, v2
 376 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 377
 378 v_cmpsx_eq_f64 vcc, -1, v[2:3]
 379 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 380
 381 v_cmpsx_eq_f64_e64 flat_scratch, v[1:2], v[2:3]
 382 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 383
 384 v_cmpsx_f_f32 vcc, -1, v2
 385 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 386
 387 v_cmpsx_f_f32_e64 flat_scratch, v1, v2
 388 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 389
 390 v_cmpsx_f_f64 vcc, -1, v[2:3]
 391 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 392
 393 v_cmpsx_f_f64_e64 flat_scratch, v[1:2], v[2:3]
 394 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 395
 396 v_cmpsx_ge_f32 vcc, -1, v2
 397 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 398
 399 v_cmpsx_ge_f32_e64 flat_scratch, v1, v2
 400 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 401
 402 v_cmpsx_ge_f64 vcc, -1, v[2:3]
 403 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 404
 405 v_cmpsx_ge_f64_e64 flat_scratch, v[1:2], v[2:3]
 406 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 407
 408 v_cmpsx_gt_f32 vcc, -1, v2
 409 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 410
 411 v_cmpsx_gt_f32_e64 flat_scratch, v1, v2
 412 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 413
 414 v_cmpsx_gt_f64 vcc, -1, v[2:3]
 415 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 416
 417 v_cmpsx_gt_f64_e64 flat_scratch, v[1:2], v[2:3]
 418 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 419
 420 v_cmpsx_le_f32 vcc, -1, v2
 421 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 422
 423 v_cmpsx_le_f32_e64 flat_scratch, v1, v2
 424 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 425
 426 v_cmpsx_le_f64 vcc, -1, v[2:3]
 427 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 428
 429 v_cmpsx_le_f64_e64 flat_scratch, v[1:2], v[2:3]
 430 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 431
 432 v_cmpsx_lg_f32 vcc, -1, v2
 433 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 434
 435 v_cmpsx_lg_f32_e64 flat_scratch, v1, v2
 436 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 437
 438 v_cmpsx_lg_f64 vcc, -1, v[2:3]
 439 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 440
 441 v_cmpsx_lg_f64_e64 flat_scratch, v[1:2], v[2:3]
 442 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 443
 444 v_cmpsx_lt_f32 vcc, -1, v2
 445 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 446
 447 v_cmpsx_lt_f32_e64 flat_scratch, v1, v2
 448 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 449
 450 v_cmpsx_lt_f64 vcc, -1, v[2:3]
 451 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 452
 453 v_cmpsx_lt_f64_e64 flat_scratch, v[1:2], v[2:3]
 454 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 455
 456 v_cmpsx_neq_f32 vcc, -1, v2
 457 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 458
 459 v_cmpsx_neq_f32_e64 flat_scratch, v1, v2
 460 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 461
 462 v_cmpsx_neq_f64 vcc, -1, v[2:3]
 463 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 464
 465 v_cmpsx_neq_f64_e64 flat_scratch, v[1:2], v[2:3]
 466 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 467
 468 v_cmpsx_nge_f32 vcc, -1, v2
 469 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 470
 471 v_cmpsx_nge_f32_e64 flat_scratch, v1, v2
 472 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 473
 474 v_cmpsx_nge_f64 vcc, -1, v[2:3]
 475 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 476
 477 v_cmpsx_nge_f64_e64 flat_scratch, v[1:2], v[2:3]
 478 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 479
 480 v_cmpsx_ngt_f32 vcc, -1, v2
 481 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 482
 483 v_cmpsx_ngt_f32_e64 flat_scratch, v1, v2
 484 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 485
 486 v_cmpsx_ngt_f64 vcc, -1, v[2:3]
 487 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 488
 489 v_cmpsx_ngt_f64_e64 flat_scratch, v[1:2], v[2:3]
 490 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 491
 492 v_cmpsx_nle_f32 vcc, -1, v2
 493 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 494
 495 v_cmpsx_nle_f32_e64 flat_scratch, v1, v2
 496 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 497
 498 v_cmpsx_nle_f64 vcc, -1, v[2:3]
 499 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 500
 501 v_cmpsx_nle_f64_e64 flat_scratch, v[1:2], v[2:3]
 502 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 503
 504 v_cmpsx_nlg_f32 vcc, -1, v2
 505 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 506
 507 v_cmpsx_nlg_f32_e64 flat_scratch, v1, v2
 508 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 509
 510 v_cmpsx_nlg_f64 vcc, -1, v[2:3]
 511 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 512
 513 v_cmpsx_nlg_f64_e64 flat_scratch, v[1:2], v[2:3]
 514 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 515
 516 v_cmpsx_nlt_f32 vcc, -1, v2
 517 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 518
 519 v_cmpsx_nlt_f32_e64 flat_scratch, v1, v2
 520 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 521
 522 v_cmpsx_nlt_f64 vcc, -1, v[2:3]
 523 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 524
 525 v_cmpsx_nlt_f64_e64 flat_scratch, v[1:2], v[2:3]
 526 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 527
 528 v_cmpsx_o_f32 vcc, -1, v2
 529 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 530
 531 v_cmpsx_o_f32_e64 flat_scratch, v1, v2
 532 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 533
 534 v_cmpsx_o_f64 vcc, -1, v[2:3]
 535 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 536
 537 v_cmpsx_o_f64_e64 flat_scratch, v[1:2], v[2:3]
 538 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 539
 540 v_cmpsx_tru_f32 vcc, -1, v2
 541 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 542
 543 v_cmpsx_tru_f32_e64 flat_scratch, v1, v2
 544 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 545
 546 v_cmpsx_tru_f64 vcc, -1, v[2:3]
 547 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 548
 549 v_cmpsx_tru_f64_e64 flat_scratch, v[1:2], v[2:3]
 550 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 551
 552 v_cmpsx_u_f32 vcc, -1, v2
 553 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 554
 555 v_cmpsx_u_f32_e64 flat_scratch, v1, v2
 556 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 557
 558 v_cmpsx_u_f64 vcc, -1, v[2:3]
 559 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 560
 561 v_cmpsx_u_f64_e64 flat_scratch, v[1:2], v[2:3]
 562 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 563
 564 v_cmpx_f_i16 vcc, -1, v2
 565 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 566
 567 v_cmpx_f_i16_e64 exec, v1, v2
 568 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 569
 570 v_cmpx_f_i16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 571 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 572
 573 v_cmpx_f_u16 vcc, -1, v2
 574 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 575
 576 v_cmpx_f_u16_e64 exec, v1, v2
 577 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 578
 579 v_cmpx_f_u16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 580 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 581
 582 v_cmpx_t_i16 vcc, -1, v2
 583 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 584
 585 v_cmpx_t_i16_e64 exec, v1, v2
 586 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 587
 588 v_cmpx_t_i16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 589 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 590
 591 v_cmpx_t_u16 vcc, -1, v2
 592 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 593
 594 v_cmpx_t_u16_e64 exec, v1, v2
 595 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 596
 597 v_cmpx_t_u16_sdwa flat_scratch, v1, v2 src0_sel:DWORD src1_sel:DWORD
 598 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 599
 600 v_cvt_pkaccum_u8_f32 v1, v2, v3
 601 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 602
 603 v_cvt_pkaccum_u8_f32_e64 v255, v1, v2
 604 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 605
 606 v_div_fixup_legacy_f16 v255, v1, v2, v3
 607 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 608
 609 v_div_fixup_legacy_f16_e64 v5, 0.5, v2, v3
 610 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 611
 612 v_dot2_f32_f16 v0, -v1, -v2, -v3
 613 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 614
 615 v_dot2_i32_i16 v0, -v1, -v2, -v3
 616 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 617
 618 v_dot2_u32_u16 v0, -v1, -v2, -v3
 619 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 620
 621 v_dot2c_f32_f16 v0, v1, v2
 622 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 623
 624 v_dot2c_f32_f16_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 625 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 626
 627 v_dot2c_f32_f16_e32 v255, v1, v2
 628 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 629
 630 v_dot2c_i32_i16 v0, v1, v2
 631 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 632
 633 v_dot2c_i32_i16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 634 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 635
 636 v_dot4_i32_i8 v0, v1, v2, v3
 637 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 638
 639 v_dot4_u32_u8 v0, v1, v2, v3
 640 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 641
 642 v_dot4c_i32_i8 v0, v1, v2
 643 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 644
 645 v_dot4c_i32_i8_dpp v255, v1, v2  quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 646 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 647
 648 v_dot4c_i32_i8_e32 v255, v1, v2
 649 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 650
 651 v_dot8_i32_i4 v0, v1, v2, v3
 652 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 653
 654 v_dot8_u32_u4 v0, v1, v2, v3
 655 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 656
 657 v_dot8c_i32_i4 v0, v1, v2
 658 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 659
 660 v_dot8c_i32_i4_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 661 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 662
 663 v_exp_legacy_f32 v1, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
 664 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 665
 666 v_exp_legacy_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 667 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 668
 669 v_exp_legacy_f32_e64 v255, v1
 670 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 671
 672 v_exp_legacy_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
 673 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 674
 675 v_fma_legacy_f16 v255, v1, v2, v3
 676 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 677
 678 v_fma_legacy_f16_e64 v5, v1, v2, v3
 679 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 680
 681 v_interp_p2_legacy_f16 v255, v2, attr0.x, v3
 682 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 683
 684 v_log_clamp_f32 v1, 0.5
 685 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 686
 687 v_log_clamp_f32_e64 v255, v1
 688 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 689
 690 v_log_legacy_f32 v1, v0 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
 691 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 692
 693 v_log_legacy_f32_dpp v255, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 694 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 695
 696 v_log_legacy_f32_e64 v255, v1
 697 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 698
 699 v_log_legacy_f32_sdwa v255, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD
 700 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 701
 702 v_lshl_b32 v255, v1, v2
 703 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 704
 705 v_lshl_b32_e64 v255, v1, v2
 706 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 707
 708 v_lshl_b64 v[254:255], v[1:2], v2
 709 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 710
 711 v_lshr_b32 v255, v1, v2
 712 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 713
 714 v_lshr_b32_e64 v255, v1, v2
 715 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 716
 717 v_lshr_b64 v[254:255], v[1:2], v2
 718 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 719
 720 v_mac_f16 v1, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
 721 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 722
 723 v_mac_f16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 724 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 725
 726 v_mac_f16_e64 v0, -4.0, flat_scratch_lo
 727 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 728
 729 v_mac_f16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
 730 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 731
 732 v_mad_f16 v255, v1, v2, v3
 733 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 734
 735 v_mad_f16_e64 v5, 0.5, v2, v3
 736 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 737
 738 v_mad_legacy_f16 v255, v1, v2, v3
 739 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 740
 741 v_mad_legacy_f16_e64 v5, 0.5, v2, v3
 742 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 743
 744 v_mad_legacy_i16 v255, v1, v2, v3
 745 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 746
 747 v_mad_legacy_i16_e64 v5, 0, v2, v3
 748 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 749
 750 v_mad_legacy_u16 v255, v1, v2, v3
 751 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 752
 753 v_mad_legacy_u16_e64 v5, 0, v2, v3
 754 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 755
 756 v_mad_mix_f32 v0, -abs(v1), v2, v3
 757 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 758
 759 v_mad_mixhi_f16 v0, -v1, abs(v2), -abs(v3)
 760 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 761
 762 v_mad_mixlo_f16 v0, abs(v1), -v2, abs(v3)
 763 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 764
 765 v_madak_f16 v0, src_lds_direct, v0, 0x1121
 766 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 767
 768 v_madmk_f16 v0, src_lds_direct, 0x1121, v0
 769 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 770
 771 v_max_legacy_f32 v255, v1, v2
 772 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 773
 774 v_max_legacy_f32_e64 v255, v1, v2
 775 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 776
 777 v_mfma_f32_16x16x16f16 a[0:3], a[0:1], a[1:2], -2.0
 778 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 779
 780 v_mfma_f32_16x16x1f32 a[0:15], a0, a1, -2.0
 781 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 782
 783 v_mfma_f32_16x16x2bf16 a[0:15], a0, a1, -2.0
 784 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 785
 786 v_mfma_f32_16x16x4f16 a[0:15], a[0:1], a[1:2], -2.0
 787 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 788
 789 v_mfma_f32_16x16x4f32 a[0:3], a0, a1, -2.0
 790 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 791
 792 v_mfma_f32_16x16x8bf16 a[0:3], a0, a1, -2.0
 793 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 794
 795 v_mfma_f32_32x32x1f32 a[0:31], 1, v1, a[1:32]
 796 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 797
 798 v_mfma_f32_32x32x2bf16 a[0:31], a0, a1, -2.0
 799 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 800
 801 v_mfma_f32_32x32x2f32 a[0:15], a0, a1, -2.0
 802 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 803
 804 v_mfma_f32_32x32x4bf16 a[0:15], a0, a1, -2.0
 805 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 806
 807 v_mfma_f32_32x32x4f16 a[0:31], a[0:1], a[1:2], -2.0
 808 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 809
 810 v_mfma_f32_32x32x8f16 a[0:15], a[0:1], a[1:2], -2.0
 811 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 812
 813 v_mfma_f32_4x4x1f32 a[0:3], a0, a1, -2.0
 814 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 815
 816 v_mfma_f32_4x4x2bf16 a[0:3], a0, a1, -2.0
 817 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 818
 819 v_mfma_f32_4x4x4f16 a[0:3], a[0:1], a[1:2], -2.0
 820 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 821
 822 v_mfma_i32_16x16x16i8 a[0:3], a0, a1, 2
 823 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 824
 825 v_mfma_i32_16x16x4i8 a[0:15], a0, a1, 2
 826 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 827
 828 v_mfma_i32_32x32x4i8 a[0:31], a0, a1, 2
 829 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 830
 831 v_mfma_i32_32x32x8i8 a[0:15], a0, a1, 2
 832 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 833
 834 v_mfma_i32_4x4x4i8 a[0:3], a0, a1, 2
 835 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 836
 837 v_min_legacy_f32 v255, v1, v2
 838 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 839
 840 v_min_legacy_f32_e64 v255, v1, v2
 841 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 842
 843 v_rcp_clamp_f32 v255, v1
 844 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 845
 846 v_rcp_clamp_f32_e64 v255, v1
 847 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 848
 849 v_rcp_clamp_f64 v[254:255], v[1:2]
 850 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 851
 852 v_rcp_clamp_f64_e64 v[254:255], v[1:2]
 853 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 854
 855 v_rcp_legacy_f32 v255, v1
 856 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 857
 858 v_rcp_legacy_f32_e64 v255, v1
 859 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 860
 861 v_rsq_clamp_f32 v255, v1
 862 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 863
 864 v_rsq_clamp_f32_e64 v255, v1
 865 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 866
 867 v_rsq_clamp_f64 v[254:255], v[1:2]
 868 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 869
 870 v_rsq_clamp_f64_e64 v[254:255], v[1:2]
 871 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 872
 873 v_rsq_legacy_f32 v255, v1
 874 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 875
 876 v_rsq_legacy_f32_e64 v255, v1
 877 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 878
 879 v_screen_partition_4se_b32 v5, -1
 880 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 881
 882 v_screen_partition_4se_b32_dpp v5, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0 bound_ctrl:0
 883 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 884
 885 v_screen_partition_4se_b32_e64 v5, -1
 886 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 887
 888 v_screen_partition_4se_b32_sdwa v5, v1 src0_sel:BYTE_0
 889 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 890
 891 v_sub_i16 v255, v1, v2
 892 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 893
 894 v_sub_i32 v1, s[0:1], v2, v3
 895 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 896
 897 v_sub_i32_e64 v255, s[12:13], v1, v2
 898 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 899
 900 v_sub_u16 v1, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
 901 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 902
 903 v_sub_u16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 904 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 905
 906 v_sub_u16_e64 v255, v1, v2
 907 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 908
 909 v_sub_u16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
 910 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 911
 912 v_sub_u32 v1, 4.0, v2
 913 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 914
 915 v_sub_u32_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 916 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 917
 918 v_sub_u32_e32 v1, s1, v3
 919 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 920
 921 v_sub_u32_e64 v255, s[12:13], v1, v2
 922 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 923
 924 v_sub_u32_sdwa v1, vcc, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
 925 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 926
 927 v_subb_co_u32 v1, vcc, v2, v3, vcc row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
 928 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 929
 930 v_subb_co_u32_dpp v255, vcc, v1, v2, vcc quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 931 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 932
 933 v_subb_co_u32_e64 v255, s[12:13], v1, v2, s[6:7]
 934 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 935
 936 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
 937 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 938
 939 v_subb_u32 v1, s[0:1], v2, v3, vcc
 940 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 941
 942 v_subb_u32_dpp v255, vcc, v1, v2, vcc quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 943 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 944
 945 v_subb_u32_e64 v255, s[12:13], v1, v2, s[6:7]
 946 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 947
 948 v_subb_u32_sdwa v1, vcc, v2, v3, vcc dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
 949 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 950
 951 v_subbrev_co_u32 v0, vcc, src_lds_direct, v0, vcc
 952 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 953
 954 v_subbrev_co_u32_dpp v255, vcc, v1, v2, vcc quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 955 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 956
 957 v_subbrev_co_u32_e64 v255, s[12:13], v1, v2, s[6:7]
 958 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 959
 960 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
 961 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 962
 963 v_subbrev_u32 v1, s[0:1], v2, v3, vcc
 964 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 965
 966 v_subbrev_u32_dpp v255, vcc, v1, v2, vcc quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 967 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 968
 969 v_subbrev_u32_e64 v255, s[12:13], v1, v2, s[6:7]
 970 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 971
 972 v_subbrev_u32_sdwa v1, vcc, v2, v3, vcc dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
 973 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 974
 975 v_subrev_i32 v1, s[0:1], v2, v3
 976 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 977
 978 v_subrev_i32_e64 v255, s[12:13], v1, v2
 979 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 980
 981 v_subrev_u16 v0, src_lds_direct, v0
 982 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 983
 984 v_subrev_u16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 985 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 986
 987 v_subrev_u16_e64 v255, v1, v2
 988 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 989
 990 v_subrev_u16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
 991 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 992
 993 v_subrev_u32 v0, src_lds_direct, v0
 994 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 995
 996 v_subrev_u32_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
 997 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 998
 999 v_subrev_u32_e32 v1, s1, v3
1000 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1001
1002 v_subrev_u32_e64 v255, s[12:13], v1, v2
1003 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1004
1005 v_subrev_u32_sdwa v1, vcc, v2, v3 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 src1_sel:BYTE_2
1006 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
1007
1008 //===----------------------------------------------------------------------===//
1009 // Unsupported e32 variants.
1010 //===----------------------------------------------------------------------===//
1011
1012 v_add_co_u32_e32 v2, vcc, s0, v2
1013 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: e32 variant of this instruction is not supported
1014
1015 v_sub_co_u32_e32 v2, vcc, s0, v2
1016 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: e32 variant of this instruction is not supported
1017
1018 v_subrev_co_u32_e32 v2, vcc, s0, v2
1019 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: e32 variant of this instruction is not supported
1020
1021 //===----------------------------------------------------------------------===//
1022 // Unsupported e64 variants.
1023 //===----------------------------------------------------------------------===//
1024
1025 v_swap_b32_e64 v1, v2
1026 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: e64 variant of this instruction is not supported
1027
1028 //===----------------------------------------------------------------------===//
1029 // Unsupported dpp variants.
1030 //===----------------------------------------------------------------------===//
1031
1032 v_add_co_u32_dpp v255, vcc, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1033 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: dpp variant of this instruction is not supported
1034
1035 v_ashrrev_i16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1036 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: dpp variant of this instruction is not supported
1037
1038 v_lshlrev_b16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1039 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: dpp variant of this instruction is not supported
1040
1041 v_lshrrev_b16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1042 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: dpp variant of this instruction is not supported
1043
1044 v_mac_legacy_f32_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1045 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: dpp variant of this instruction is not supported
1046
1047 v_max_i16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1048 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: dpp variant of this instruction is not supported
1049
1050 v_max_u16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1051 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: dpp variant of this instruction is not supported
1052
1053 v_min_i16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1054 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: dpp variant of this instruction is not supported
1055
1056 v_min_u16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1057 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: dpp variant of this instruction is not supported
1058
1059 v_mul_lo_u16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1060 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: dpp variant of this instruction is not supported
1061
1062 v_sub_co_u32_dpp v255, vcc, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1063 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: dpp variant of this instruction is not supported
1064
1065 v_subrev_co_u32_dpp v255, vcc, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
1066 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: dpp variant of this instruction is not supported
1067
1068 //===----------------------------------------------------------------------===//
1069 // Unsupported sdwa variants.
1070 //===----------------------------------------------------------------------===//
1071
1072 v_add_co_u32_sdwa v0, v0, v4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:BYTE_0
1073 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: sdwa variant of this instruction is not supported
1074
1075 v_ashrrev_i16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1076 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: sdwa variant of this instruction is not supported
1077
1078 v_lshlrev_b16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1079 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: sdwa variant of this instruction is not supported
1080
1081 v_lshrrev_b16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1082 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: sdwa variant of this instruction is not supported
1083
1084 v_mac_f32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1085 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: sdwa variant of this instruction is not supported
1086
1087 v_mac_legacy_f32_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1088 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: sdwa variant of this instruction is not supported
1089
1090 v_max_i16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1091 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: sdwa variant of this instruction is not supported
1092
1093 v_max_u16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1094 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: sdwa variant of this instruction is not supported
1095
1096 v_min_i16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1097 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: sdwa variant of this instruction is not supported
1098
1099 v_min_u16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1100 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: sdwa variant of this instruction is not supported
1101
1102 v_mul_lo_u16_sdwa v255, v1, v2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:DWORD
1103 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: sdwa variant of this instruction is not supported
1104
1105 v_sub_co_u32_sdwa v0, v0, v4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:BYTE_0
1106 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: sdwa variant of this instruction is not supported
1107
1108 v_subrev_co_u32_sdwa v0, v0, v4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:DWORD src1_sel:BYTE_0
1109 // CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: sdwa variant of this instruction is not supported