xref: /llvm-project/llvm/test/CodeGen/AMDGPU/hsa-metadata-kernel-code-props.ll (revision b1bcb7ca460fcd317bbc8309e14c8761bf8394e0)
1; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -enable-misched=0 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefixes=CHECK,GFX700,WAVE64 %s
2; RUN: llc -mattr=-xnack -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 -enable-misched=0 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefixes=CHECK,GFX803,WAVE64 %s
3; RUN: llc -mattr=-xnack -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -enable-misched=0 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefixes=CHECK,GFX900,WAVE64 %s
4; RUN: llc -mattr=-xnack -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1010 -enable-misched=0 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefixes=CHECK,GFX1010,WAVE32 %s
5
6@var = addrspace(1) global float 0.0
7
8; CHECK: ---
9; CHECK:  amdhsa.kernels:
10
11; CHECK:   - .args:
12; CHECK:     .group_segment_fixed_size: 0
13; CHECK:     .kernarg_segment_align: 8
14; CHECK:     .kernarg_segment_size: 24
15; CHECK:     .max_flat_workgroup_size: 1024
16; CHECK:     .name:           test
17; CHECK:     .private_segment_fixed_size: 0
18; CHECK:     .sgpr_count:     10
19; CHECK:     .symbol:         test.kd
20; CHECK:     .vgpr_count:     {{3|6}}
21; WAVE64:    .wavefront_size: 64
22; WAVE32:    .wavefront_size: 32
23define amdgpu_kernel void @test(
24    ptr addrspace(1) %r,
25    ptr addrspace(1) %a,
26    ptr addrspace(1) %b) "amdgpu-no-implicitarg-ptr" {
27entry:
28  %a.val = load half, ptr addrspace(1) %a
29  %b.val = load half, ptr addrspace(1) %b
30  %r.val = fadd half %a.val, %b.val
31  store half %r.val, ptr addrspace(1) %r
32  ret void
33}
34
35; CHECK:   - .args:
36; CHECK:     .max_flat_workgroup_size: 256
37define amdgpu_kernel void @test_max_flat_workgroup_size(
38    ptr addrspace(1) %r,
39    ptr addrspace(1) %a,
40    ptr addrspace(1) %b) #2 {
41entry:
42  %a.val = load half, ptr addrspace(1) %a
43  %b.val = load half, ptr addrspace(1) %b
44  %r.val = fadd half %a.val, %b.val
45  store half %r.val, ptr addrspace(1) %r
46  ret void
47}
48
49; CHECK:   .name:       num_spilled_sgprs
50; GFX700:   .sgpr_spill_count: 10
51; GFX803:   .sgpr_spill_count: 10
52; GFX900:   .sgpr_spill_count: 62
53; GFX1010:  .sgpr_spill_count: 60
54; CHECK:   .symbol:     num_spilled_sgprs.kd
55define amdgpu_kernel void @num_spilled_sgprs(
56    ptr addrspace(1) %out0, ptr addrspace(1) %out1, [8 x i32],
57    ptr addrspace(1) %out2, ptr addrspace(1) %out3, [8 x i32],
58    ptr addrspace(1) %out4, ptr addrspace(1) %out5, [8 x i32],
59    ptr addrspace(1) %out6, ptr addrspace(1) %out7, [8 x i32],
60    ptr addrspace(1) %out8, ptr addrspace(1) %out9, [8 x i32],
61    ptr addrspace(1) %outa, ptr addrspace(1) %outb, [8 x i32],
62    ptr addrspace(1) %outc, ptr addrspace(1) %outd, [8 x i32],
63    ptr addrspace(1) %oute, ptr addrspace(1) %outf, [8 x i32],
64    ptr addrspace(1) %outg, ptr addrspace(1) %outh, [8 x i32],
65    ptr addrspace(1) %outi, ptr addrspace(1) %outj, [8 x i32],
66    ptr addrspace(1) %outk, ptr addrspace(1) %outl, [8 x i32],
67    ptr addrspace(1) %outm, ptr addrspace(1) %outn, [8 x i32],
68    i32 %in0, i32 %in1, i32 %in2, i32 %in3, [8 x i32],
69    i32 %in4, i32 %in5, i32 %in6, i32 %in7, [8 x i32],
70    i32 %in8, i32 %in9, i32 %ina, i32 %inb, [8 x i32],
71    i32 %inc, i32 %ind, i32 %ine, i32 %inf, i32 %ing, i32 %inh,
72    i32 %ini, i32 %inj, i32 %ink) #0 {
73entry:
74  store volatile i32 %in0, ptr addrspace(1) %out0
75  store volatile i32 %in1, ptr addrspace(1) %out1
76  store volatile i32 %in2, ptr addrspace(1) %out2
77  store volatile i32 %in3, ptr addrspace(1) %out3
78  store volatile i32 %in4, ptr addrspace(1) %out4
79  store volatile i32 %in5, ptr addrspace(1) %out5
80  store volatile i32 %in6, ptr addrspace(1) %out6
81  store volatile i32 %in7, ptr addrspace(1) %out7
82  store volatile i32 %in8, ptr addrspace(1) %out8
83  store volatile i32 %in9, ptr addrspace(1) %out9
84  store volatile i32 %ina, ptr addrspace(1) %outa
85  store volatile i32 %inb, ptr addrspace(1) %outb
86  store volatile i32 %inc, ptr addrspace(1) %outc
87  store volatile i32 %ind, ptr addrspace(1) %outd
88  store volatile i32 %ine, ptr addrspace(1) %oute
89  store volatile i32 %inf, ptr addrspace(1) %outf
90  store volatile i32 %ing, ptr addrspace(1) %outg
91  store volatile i32 %inh, ptr addrspace(1) %outh
92  store volatile i32 %ini, ptr addrspace(1) %outi
93  store volatile i32 %inj, ptr addrspace(1) %outj
94  store volatile i32 %ink, ptr addrspace(1) %outk
95  ret void
96}
97
98; CHECK:   .name:       num_spilled_vgprs
99; CHECK:   .symbol:     num_spilled_vgprs.kd
100; CHECK:   .vgpr_spill_count: {{13|14}}
101define amdgpu_kernel void @num_spilled_vgprs() #1 {
102  %val0 = load volatile float, ptr addrspace(1) @var
103  %val1 = load volatile float, ptr addrspace(1) @var
104  %val2 = load volatile float, ptr addrspace(1) @var
105  %val3 = load volatile float, ptr addrspace(1) @var
106  %val4 = load volatile float, ptr addrspace(1) @var
107  %val5 = load volatile float, ptr addrspace(1) @var
108  %val6 = load volatile float, ptr addrspace(1) @var
109  %val7 = load volatile float, ptr addrspace(1) @var
110  %val8 = load volatile float, ptr addrspace(1) @var
111  %val9 = load volatile float, ptr addrspace(1) @var
112  %val10 = load volatile float, ptr addrspace(1) @var
113  %val11 = load volatile float, ptr addrspace(1) @var
114  %val12 = load volatile float, ptr addrspace(1) @var
115  %val13 = load volatile float, ptr addrspace(1) @var
116  %val14 = load volatile float, ptr addrspace(1) @var
117  %val15 = load volatile float, ptr addrspace(1) @var
118  %val16 = load volatile float, ptr addrspace(1) @var
119  %val17 = load volatile float, ptr addrspace(1) @var
120  %val18 = load volatile float, ptr addrspace(1) @var
121  %val19 = load volatile float, ptr addrspace(1) @var
122  %val20 = load volatile float, ptr addrspace(1) @var
123  %val21 = load volatile float, ptr addrspace(1) @var
124  %val22 = load volatile float, ptr addrspace(1) @var
125  %val23 = load volatile float, ptr addrspace(1) @var
126  %val24 = load volatile float, ptr addrspace(1) @var
127  %val25 = load volatile float, ptr addrspace(1) @var
128  %val26 = load volatile float, ptr addrspace(1) @var
129  %val27 = load volatile float, ptr addrspace(1) @var
130  %val28 = load volatile float, ptr addrspace(1) @var
131  %val29 = load volatile float, ptr addrspace(1) @var
132  %val30 = load volatile float, ptr addrspace(1) @var
133
134  store volatile float %val0, ptr addrspace(1) @var
135  store volatile float %val1, ptr addrspace(1) @var
136  store volatile float %val2, ptr addrspace(1) @var
137  store volatile float %val3, ptr addrspace(1) @var
138  store volatile float %val4, ptr addrspace(1) @var
139  store volatile float %val5, ptr addrspace(1) @var
140  store volatile float %val6, ptr addrspace(1) @var
141  store volatile float %val7, ptr addrspace(1) @var
142  store volatile float %val8, ptr addrspace(1) @var
143  store volatile float %val9, ptr addrspace(1) @var
144  store volatile float %val10, ptr addrspace(1) @var
145  store volatile float %val11, ptr addrspace(1) @var
146  store volatile float %val12, ptr addrspace(1) @var
147  store volatile float %val13, ptr addrspace(1) @var
148  store volatile float %val14, ptr addrspace(1) @var
149  store volatile float %val15, ptr addrspace(1) @var
150  store volatile float %val16, ptr addrspace(1) @var
151  store volatile float %val17, ptr addrspace(1) @var
152  store volatile float %val18, ptr addrspace(1) @var
153  store volatile float %val19, ptr addrspace(1) @var
154  store volatile float %val20, ptr addrspace(1) @var
155  store volatile float %val21, ptr addrspace(1) @var
156  store volatile float %val22, ptr addrspace(1) @var
157  store volatile float %val23, ptr addrspace(1) @var
158  store volatile float %val24, ptr addrspace(1) @var
159  store volatile float %val25, ptr addrspace(1) @var
160  store volatile float %val26, ptr addrspace(1) @var
161  store volatile float %val27, ptr addrspace(1) @var
162  store volatile float %val28, ptr addrspace(1) @var
163  store volatile float %val29, ptr addrspace(1) @var
164  store volatile float %val30, ptr addrspace(1) @var
165
166  ret void
167}
168
169; CHECK:  amdhsa.version:
170; CHECK-NEXT: - 1
171; CHECK-NEXT: - 1
172
173attributes #0 = { "amdgpu-num-sgpr"="20" }
174attributes #1 = { "amdgpu-num-vgpr"="20" }
175attributes #2 = { "amdgpu-flat-work-group-size"="1,256" }
176
177!llvm.module.flags = !{!0}
178!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}
179