llvm/test/MC/AMDGPU/mai-gfx950-err.s

// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx950 %s 2>&1 | FileCheck --implicit-check-not=error: %s

v_mfma_ld_scale_b32 v0, 65
// CHECK: :[[@LINE-1]]:25: error: literal operands are not supported

v_mfma_ld_scale_b32 65, v0
// CHECK: :[[@LINE-1]]:21: error: literal operands are not supported

v_mfma_ld_scale_b32 65, 65
// CHECK: :[[@LINE-1]]:25: error: literal operands are not supported

v_mfma_ld_scale_b32 s0, s1
// CHECK: :[[@LINE-1]]:25: error: invalid operand (violates constant bus restrictions)

v_mfma_ld_scale_b32 v0, v0 clamp
// CHECK: :[[@LINE-1]]:28: error: invalid operand for instruction

v_mfma_ld_scale_b32 v0, v0 neg_lo:[0,1]
// CHECK: :[[@LINE-1]]:28: error: not a valid operand

v_mfma_ld_scale_b32 v0, v0 neg_lo:[1,1]
// CHECK: :[[@LINE-1]]:28: error: not a valid operand

v_mfma_ld_scale_b32 v0, v0 neg_hi:[1,1]
// CHECK: :[[@LINE-1]]:28: error: not a valid operand

v_mfma_ld_scale_b32 v0, v0 neg_hi:[0,1]
// CHECK: :[[@LINE-1]]:28: error: not a valid operand

v_mfma_ld_scale_b32 v0, v0 neg_lo:[0,1] neg_hi:[0,1]
// CHECK: :[[@LINE-1]]:28: error: not a valid operand


v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] cbsz:2
// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 2

v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] blgp:2
// CHECK: :[[@LINE-1]]:46: error: wrong register tuple size for blgp value 2

v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] cbsz:2 blgp:2
// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 2
// CHECK: :[[@LINE-2]]:46: error: wrong register tuple size for blgp value 2


v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] cbsz:2
// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 2

v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] blgp:2
// CHECK: :[[@LINE-1]]:46: error: wrong register tuple size for blgp value 2

v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] cbsz:2 blgp:2
// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 2
// CHECK: :[[@LINE-2]]:46: error: wrong register tuple size for blgp value 2

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] v20, v21 cbsz:2
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] v20, v21 cbsz:2
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] v20, v21 blgp:2
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 2

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] v20, v21 blgp:2
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 2

v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15] cbsz:2
// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 2

v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15] blgp:2
// CHECK: :[[@LINE-1]]:47: error: wrong register tuple size for blgp value 2

v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[0:15] cbsz:2
// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 2

v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[0:15] blgp:2
// CHECK: :[[@LINE-1]]:47: error: wrong register tuple size for blgp value 2

v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15] v32, v33 cbsz:2
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2

v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[0:15] v32, v33 cbsz:2
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2



v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 0

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21 cbsz:1
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 1

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21 cbsz:2
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21 cbsz:3
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 3


v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[0:3] v20, v21
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 0

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[0:3] v20, v21 blgp:1
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 1

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[0:3] v20, v21 blgp:2
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 2

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[0:3] v20, v21 blgp:3
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 3

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[0:3] v20, v21
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 0

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[0:3] v20, v21 cbsz:1
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 1

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[0:3] v20, v21 cbsz:2
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[0:3] v20, v21 cbsz:3
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 3

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[0:3] v20, v21
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 0

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[0:3] v20, v21 blgp:1
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 1

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[0:3] v20, v21 blgp:2
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 2

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[0:3] v20, v21 blgp:3
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 3

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21 cbsz:3
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 3

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] v20, v21 cbsz:3
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 3

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[12:19], v[4:7], v[0:3] v20, v21 blgp:3
// CHECK: :[[@LINE-1]]:53: error: wrong register tuple size for blgp value 3

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[12:19], v[4:11], v[0:3] v20, v21 blgp:3
// CHECK: :[[@LINE-1]]:53: error: wrong register tuple size for blgp value 3

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[0:3] v20, v21 cbsz:4
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 4

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] v20, v21 cbsz:4
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 4

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[12:19], v[4:9], v[0:3] v20, v21 blgp:4
// CHECK: :[[@LINE-1]]:53: error: wrong register tuple size for blgp value 4

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[12:19], v[4:11], v[0:3] v20, v21 blgp:4
// CHECK: :[[@LINE-1]]:53: error: wrong register tuple size for blgp value 4