253 lines · plain
1// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx950 %s 2>&1 | FileCheck -check-prefix=ERR %s2// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx950 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck -check-prefix=W32-ERR %s3 4//===----------------------------------------------------------------------===//5// v_mfma_f32_32x32x4_xf326//===----------------------------------------------------------------------===//7 8v_mfma_f32_32x32x4_xf32 a[0:3], v[2:3], v[4:5], a[2:5]9// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU10 11v_mfma_f32_32x32x4_xf32 v[0:3], v[0:3], v[0:3], v[0:3]12// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU13 14v_mfma_f32_32x32x4_xf32 v[0:3], v[0:3], v[0:3], v[0:3]15// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU16 17v_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3]18// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU19 20v_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3]21// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU22 23v_mfma_f32_32x32x4_xf32 v[0:3], a[0:3], v[0:3], 1.024// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU25 26v_mfma_f32_32x32x4_xf32 a[0:3], v[0:3], a[0:3], 1.027// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU28 29v_mfma_f32_32x32x4_xf32 v[0:3], v[0:3], v[0:3], v[0:3] blgp:530// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU31 32v_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] blgp:133// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU34 35v_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:336// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU37 38v_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] abid:139// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU40 41v_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:142// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU43 44v_mfma_f32_32x32x4_xf32 a[0:3], v[0:3], v[0:3], a[4:7]45// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU46 47v_mfma_f32_32x32x4_xf32 v[0:3], a[0:3], a[0:3], v[4:7]48// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU49 50v_mfma_f32_32x32x4_xf32 a[0:3], v[2:3], v[4:5], a[2:5]51// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU52 53v_mfma_f32_32x32x4_xf32 v[0:3], v[0:3], v[0:3], v[0:3]54// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU55 56v_mfma_f32_32x32x4_xf32 v[0:3], v[0:3], v[0:3], v[0:3]57// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU58 59v_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3]60// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU61 62v_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3]63// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU64 65v_mfma_f32_32x32x4_xf32 v[0:3], a[0:3], v[0:3], 1.066// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU67 68v_mfma_f32_32x32x4_xf32 a[0:3], v[0:3], a[0:3], 1.069// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU70 71v_mfma_f32_32x32x4_xf32 v[0:3], v[0:3], v[0:3], v[0:3] blgp:572// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU73 74v_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] blgp:175// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU76 77v_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:378// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU79 80v_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] abid:181// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU82 83v_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:184// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU85 86v_mfma_f32_32x32x4_xf32 a[0:3], v[0:3], v[0:3], a[4:7]87// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU88 89v_mfma_f32_32x32x4_xf32 v[0:3], a[0:3], a[0:3], v[4:7]90// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU91 92 93//===----------------------------------------------------------------------===//94// v_mfma_f32_16x16x8_xf3295//===----------------------------------------------------------------------===//96 97v_mfma_f32_16x16x8_xf32 a[0:3], v[2:3], v[4:5], a[2:5]98// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU99 100v_mfma_f32_16x16x8_xf32 v[0:3], v[0:3], v[0:3], v[0:3]101// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU102 103v_mfma_f32_16x16x8_xf32 v[0:3], v[0:3], v[0:3], v[0:3]104// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU105 106v_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3]107// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU108 109v_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3]110// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU111 112v_mfma_f32_16x16x8_xf32 v[0:3], a[0:3], v[0:3], 1.0113// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU114 115v_mfma_f32_16x16x8_xf32 a[0:3], v[0:3], a[0:3], 1.0116// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU117 118v_mfma_f32_16x16x8_xf32 v[0:3], v[0:3], v[0:3], v[0:3] blgp:5119// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU120 121v_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] blgp:1122// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU123 124v_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3125// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU126 127v_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] abid:1128// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU129 130v_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:1131// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU132 133v_mfma_f32_16x16x8_xf32 a[0:3], v[0:3], v[0:3], a[4:7]134// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU135 136v_mfma_f32_16x16x8_xf32 v[0:3], a[0:3], a[0:3], v[4:7]137// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU138 139 140v_mfma_f32_16x16x8_xf32 a[0:3], v[2:3], v[4:5], a[2:5]141// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU142 143v_mfma_f32_16x16x8_xf32 v[0:3], v[0:3], v[0:3], v[0:3]144// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU145 146v_mfma_f32_16x16x8_xf32 v[0:3], v[0:3], v[0:3], v[0:3]147// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU148 149v_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3]150// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU151 152v_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3]153// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU154 155v_mfma_f32_16x16x8_xf32 v[0:3], a[0:3], v[0:3], 1.0156// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU157 158v_mfma_f32_16x16x8_xf32 a[0:3], v[0:3], a[0:3], 1.0159// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU160 161v_mfma_f32_16x16x8_xf32 v[0:3], v[0:3], v[0:3], v[0:3] blgp:5162// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU163 164v_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] blgp:1165// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU166 167v_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3168// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU169 170v_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] abid:1171// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU172 173v_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:1174// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU175 176v_mfma_f32_16x16x8_xf32 a[0:3], v[0:3], v[0:3], a[4:7]177// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU178 179v_mfma_f32_16x16x8_xf32 v[0:3], a[0:3], a[0:3], v[4:7]180// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU181 182//===----------------------------------------------------------------------===//183// ds_read_b64_tr_b4184//===----------------------------------------------------------------------===//185ds_read_b64_tr_b4 v[1:2], v0186// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction187// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU188 189ds_read_b64_tr_b4 v1, v0190// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction191// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU192 193ds_read_b64_tr_b4 v[0:1], s0194// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction195// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU196 197ds_read_b64_tr_b4 v[2:3], v2 offset:-64198// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: expected a 16-bit unsigned offset199// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU200 201//===----------------------------------------------------------------------===//202//ds_read_b64_tr_b8203//===----------------------------------------------------------------------===//204ds_read_b64_tr_b8 v[1:2], v0205// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction206// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU207 208ds_read_b64_tr_b8 v1, v0209// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction210// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU211 212ds_read_b64_tr_b8 v[0:1], s0213// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction214// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU215 216ds_read_b64_tr_b8 v[2:3], v2 offset:-64217// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: expected a 16-bit unsigned offset218// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU219 220//===----------------------------------------------------------------------===//221// ds_read_b64_tr_b16222//===----------------------------------------------------------------------===//223ds_read_b64_tr_b16 v[1:2], v0224// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction225// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU226 227ds_read_b64_tr_b16 v1, v0228// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction229// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU230 231ds_read_b64_tr_b16 v[0:1], s0232// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction233// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU234 235ds_read_b64_tr_b16 v[2:3], v2 offset:-64236// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: expected a 16-bit unsigned offset237// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU238 239//===----------------------------------------------------------------------===//240// ds_read_b96_tr_b6241//===----------------------------------------------------------------------===//242ds_read_b96_tr_b6 v1, v0243// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction244// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU245 246ds_read_b96_tr_b6 v[0:3], s0247// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction248// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU249 250ds_read_b96_tr_b6 v[2:4], v2 offset:-64251// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: expected a 16-bit unsigned offset252// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU253