1// RUN: llvm-mc -arch=amdgcn -mcpu=gfx90a %s 2>&1 | FileCheck -check-prefixes=GFX90A %s
2// RUN: llvm-mc -arch=amdgcn -mcpu=gfx908 %s 2>&1 | FileCheck -check-prefixes=GFX908 %s
3// Based on sym_kernel_scope.s
4
5.byte .kernel.agpr_count
6// CHECK: .byte 0
7.byte .kernel.vgpr_count
8// CHECK: .byte 0
9
10    v_accvgpr_write_b32 a0, v6
11    v_accvgpr_read_b32 v3, a3
12    s_endpgm
13.byte .kernel.agpr_count
14// GFX90A: .byte 4
15// GFX908: .byte 4
16.byte .kernel.vgpr_count
17// GFX90A: .byte 12
18// GFX908: .byte 7
19
20.amdgpu_hsa_kernel K1
21K1:
22.byte .kernel.agpr_count
23// CHECK: .byte 0
24.byte .kernel.vgpr_count
25// CHECK: .byte 0
26    v_accvgpr_write_b32 a44, v6
27    s_endpgm
28.byte .kernel.agpr_count
29// GFX90A: .byte 45
30// GFX908: .byte 45
31.byte .kernel.vgpr_count
32// GFX90A: .byte 53
33// GFX908: .byte 45
34
35.amdgpu_hsa_kernel K2
36.byte .kernel.agpr_count
37// CHECK: .byte 0
38.byte .kernel.vgpr_count
39// CHECK: .byte 0
40K2:
41    v_mfma_f32_4x4x1f32 a[0:3], v1, v0, a[0:3] cbsz:1 abid:2 blgp:3
42    s_endpgm
43.byte .kernel.agpr_count
44// GFX90A: .byte 4
45// GFX908: .byte 4
46.byte .kernel.vgpr_count
47// GFX90A: .byte 8
48// GFX908: .byte 4
49
50.text
51.amdgpu_hsa_kernel K3
52K3:
53    v_accvgpr_read_b32 v[0], a0
54    v_mfma_f32_16x16x1f32 a[0:15], v1, v0, a[0:15] cbsz:1 abid:2 blgp:3
55    s_endpgm
56
57.byte .kernel.agpr_count
58// GFX90A: .byte 16
59// GFX908: .byte 16
60.byte .kernel.vgpr_count
61// GFX90A: .byte 20
62// GFX908: .byte 16
63