209 lines · plain
1// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx950 %s 2>&1 | FileCheck --implicit-check-not=error: %s2 3v_mfma_ld_scale_b32 v0, 654// CHECK: :[[@LINE-1]]:25: error: literal operands are not supported5 6v_mfma_ld_scale_b32 65, v07// CHECK: :[[@LINE-1]]:21: error: literal operands are not supported8 9v_mfma_ld_scale_b32 65, 6510// CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: literal operands are not supported11// CHECK-NEXT:{{^}}v_mfma_ld_scale_b32 65, 6512// CHECK-NEXT:{{^}} ^13 14v_mfma_ld_scale_b32 s0, s115// CHECK: :[[@LINE-1]]:25: error: invalid operand (violates constant bus restrictions)16 17v_mfma_ld_scale_b32 v0, v0 clamp18// CHECK: :[[@LINE-1]]:28: error: invalid operand for instruction19 20v_mfma_ld_scale_b32 v0, v0 neg_lo:[0,1]21// CHECK: :[[@LINE-1]]:28: error: not a valid operand22 23v_mfma_ld_scale_b32 v0, v0 neg_lo:[1,1]24// CHECK: :[[@LINE-1]]:28: error: not a valid operand25 26v_mfma_ld_scale_b32 v0, v0 neg_hi:[1,1]27// CHECK: :[[@LINE-1]]:28: error: not a valid operand28 29v_mfma_ld_scale_b32 v0, v0 neg_hi:[0,1]30// CHECK: :[[@LINE-1]]:28: error: not a valid operand31 32v_mfma_ld_scale_b32 v0, v0 neg_lo:[0,1] neg_hi:[0,1]33// CHECK: :[[@LINE-1]]:28: error: not a valid operand34 35 36v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] cbsz:237// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 238 39v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] blgp:240// CHECK: :[[@LINE-1]]:46: error: wrong register tuple size for blgp value 241 42v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] cbsz:2 blgp:243// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 244// CHECK: :[[@LINE-2]]:46: error: wrong register tuple size for blgp value 245 46 47v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] cbsz:248// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 249 50v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] blgp:251// CHECK: :[[@LINE-1]]:46: error: wrong register tuple size for blgp value 252 53v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] cbsz:2 blgp:254// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 255// CHECK: :[[@LINE-2]]:46: error: wrong register tuple size for blgp value 256 57v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] v20, v21 cbsz:258// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 259 60v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] v20, v21 cbsz:261// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 262 63v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] v20, v21 blgp:264// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 265 66v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] v20, v21 blgp:267// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 268 69v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15] cbsz:270// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 271 72v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15] blgp:273// CHECK: :[[@LINE-1]]:47: error: wrong register tuple size for blgp value 274 75v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[0:15] cbsz:276// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 277 78v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[0:15] blgp:279// CHECK: :[[@LINE-1]]:47: error: wrong register tuple size for blgp value 280 81v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15] v32, v33 cbsz:282// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 283 84v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[0:15] v32, v33 cbsz:285// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 286 87 88 89v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v2190// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 091 92v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21 cbsz:193// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 194 95v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21 cbsz:296// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 297 98v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21 cbsz:399// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 3100 101 102v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[0:3] v20, v21103// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 0104 105v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[0:3] v20, v21 blgp:1106// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 1107 108v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[0:3] v20, v21 blgp:2109// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 2110 111v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[0:3] v20, v21 blgp:3112// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 3113 114v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[0:3] v20, v21115// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 0116 117v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[0:3] v20, v21 cbsz:1118// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 1119 120v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[0:3] v20, v21 cbsz:2121// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2122 123v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[0:3] v20, v21 cbsz:3124// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 3125 126v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[0:3] v20, v21127// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 0128 129v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[0:3] v20, v21 blgp:1130// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 1131 132v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[0:3] v20, v21 blgp:2133// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 2134 135v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[0:3] v20, v21 blgp:3136// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 3137 138v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21 cbsz:3139// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 3140 141v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] v20, v21 cbsz:3142// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 3143 144v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[12:19], v[4:7], v[0:3] v20, v21 blgp:3145// CHECK: :[[@LINE-1]]:53: error: wrong register tuple size for blgp value 3146 147v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[12:19], v[4:11], v[0:3] v20, v21 blgp:3148// CHECK: :[[@LINE-1]]:53: error: wrong register tuple size for blgp value 3149 150v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[0:3] v20, v21 cbsz:4151// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 4152 153v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] v20, v21 cbsz:4154// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 4155 156v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[12:19], v[4:9], v[0:3] v20, v21 blgp:4157// CHECK: :[[@LINE-1]]:53: error: wrong register tuple size for blgp value 4158 159v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[12:19], v[4:11], v[0:3] v20, v21 blgp:4160// CHECK: :[[@LINE-1]]:53: error: wrong register tuple size for blgp value 4161 162 163// Workaround a hardware bug to disallow sgpr/inline constants as scale operands164 165v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v44, s24166// CHECK: :[[@LINE-1]]:77: error: invalid operand for instruction167 168v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], s24, v44169// CHECK: :[[@LINE-1]]:72: error: invalid operand for instruction170 171v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], m0, v24172// CHECK: :[[@LINE-1]]:72: error: invalid operand for instruction173 174v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], vcc_lo, v24175// CHECK: :[[@LINE-1]]:72: error: invalid operand for instruction176 177v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], 9, v24178// CHECK: :[[@LINE-1]]:72: error: invalid operand for instruction179 180v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, 9181// CHECK: :[[@LINE-1]]:77: error: invalid operand for instruction182 183v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], 33, v24184// CHECK: :[[@LINE-1]]:72: error: invalid operand for instruction185 186v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], 4.0, v24187// CHECK: :[[@LINE-1]]:72: error: invalid operand for instruction188 189v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], v24, 4.0190// CHECK: :[[@LINE-1]]:77: error: invalid operand for instruction191 192v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], -4.0, v24193// CHECK: :[[@LINE-1]]:72: error: invalid operand for instruction194 195v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[20:23], 0.15915494, v24196// CHECK: :[[@LINE-1]]:72: error: invalid operand for instruction197 198v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], 16, v49199// CHECK: :[[@LINE-1]]:73: error: invalid operand for instruction200 201v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], v48, -4.0202// CHECK: :[[@LINE-1]]:78: error: invalid operand for instruction203 204v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], 4.0, v24205// CHECK: :[[@LINE-1]]:73: error: invalid operand for instruction206 207v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[32:47], 0.15915494, v24208// CHECK: :[[@LINE-1]]:73: error: invalid operand for instruction209