llvm/test/MC/AMDGPU/gfx9_unsupported.s

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