Skip to content

Commit ae4a81e

Browse files
abhinavgabaabhinavgaba
authored
[NFC][OpenMP] Add tests for mapping pointers and their dereferences. (#146934)
The output of the compile-and-run tests is incorrect. These will be used for reference in future commits that resolve the issues. Also updated the existing clang LIT test, target_map_both_pointer_pointee_codegen.cpp, with more constructs and fewer CHECKs (through more update_cc_test_checks filters).
1 parent 1830b87 commit ae4a81e

12 files changed

+1874
-156
lines changed

‎clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp

Lines changed: 237 additions & 156 deletions
Large diffs are not rendered by default.

‎clang/test/OpenMP/target_map_both_pointer_pointee_codegen_global.cpp

Lines changed: 249 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,249 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --check-globals --filter-out-after "getelem.*kernel" --filter-out "= alloca.*" --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _ --global-value-regex "\.offload_.*" --global-hex-value-regex ".offload_maptypes.*"
2+
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
3+
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
4+
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
5+
6+
// expected-no-diagnostics
7+
#ifndef HEADER
8+
#define HEADER
9+
10+
int *ptr;
11+
12+
void f1() {
13+
// &ptr, &ptr, sizeof(ptr), TO | PARAM
14+
#pragma omp target map(ptr)
15+
ptr[1] = 5;
16+
}
17+
18+
void f2() {
19+
// &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ
20+
#pragma omp target map(ptr[2])
21+
ptr[1] = 6;
22+
}
23+
24+
void f3() {
25+
// &ptr, &ptr[0], sizeof(ptr[0:2]), TO | FROM | PARAM | PTR_AND_OBJ
26+
#pragma omp target map(ptr, ptr[0:2])
27+
ptr[1] = 7;
28+
}
29+
30+
void f4() {
31+
// &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ
32+
#pragma omp target map(ptr, ptr[2])
33+
ptr[2] = 8;
34+
}
35+
36+
void f5() {
37+
// &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ
38+
#pragma omp target map(ptr[2], ptr)
39+
ptr[2] = 9;
40+
}
41+
42+
void f6() {
43+
// &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ
44+
// TODO: PARAM should not be needed here.
45+
#pragma omp target data map(ptr, ptr[2])
46+
ptr[2] = 10;
47+
}
48+
49+
void f7() {
50+
// &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ
51+
// TODO: PARAM should not be needed here.
52+
#pragma omp target data map(ptr[2], ptr)
53+
ptr[2] = 11;
54+
}
55+
#endif
56+
//.
57+
// CHECK: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 8]
58+
// CHECK: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x23]]]
59+
// CHECK: @.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 4]
60+
// CHECK: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]]
61+
// CHECK: @.offload_sizes.3 = private unnamed_addr constant [1 x i64] [i64 8]
62+
// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]]
63+
// CHECK: @.offload_sizes.5 = private unnamed_addr constant [1 x i64] [i64 4]
64+
// CHECK: @.offload_maptypes.6 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]]
65+
// CHECK: @.offload_sizes.7 = private unnamed_addr constant [1 x i64] [i64 4]
66+
// CHECK: @.offload_maptypes.8 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]]
67+
// CHECK: @.offload_sizes.9 = private unnamed_addr constant [1 x i64] [i64 4]
68+
// CHECK: @.offload_maptypes.10 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]]
69+
// CHECK: @.offload_sizes.11 = private unnamed_addr constant [1 x i64] [i64 4]
70+
// CHECK: @.offload_maptypes.12 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]]
71+
//.
72+
// CHECK-LABEL: define {{[^@]+}}@_Z2f1v
73+
// CHECK-SAME: () #[[ATTR0:[0-9]+]] {
74+
// CHECK: entry:
75+
// CHECK: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
76+
// CHECK: store ptr @ptr, ptr [[TMP0]], align 8
77+
// CHECK: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
78+
// CHECK: store ptr @ptr, ptr [[TMP1]], align 8
79+
// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
80+
// CHECK: store ptr null, ptr [[TMP2]], align 8
81+
// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
82+
// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
83+
// CHECK: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0
84+
//
85+
//
86+
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f1v_l14
87+
// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[PTR:%.*]]) #[[ATTR1:[0-9]+]] {
88+
// CHECK: entry:
89+
// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8
90+
// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8, !nonnull [[META13:![0-9]+]], !align [[META14:![0-9]+]]
91+
// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8
92+
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 1
93+
// CHECK: store i32 5, ptr [[ARRAYIDX]], align 4
94+
// CHECK: ret void
95+
//
96+
//
97+
// CHECK-LABEL: define {{[^@]+}}@_Z2f2v
98+
// CHECK-SAME: () #[[ATTR0]] {
99+
// CHECK: entry:
100+
// CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8
101+
// CHECK: [[TMP1:%.*]] = load ptr, ptr @ptr, align 8
102+
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 2
103+
// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
104+
// CHECK: store ptr @ptr, ptr [[TMP2]], align 8
105+
// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
106+
// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8
107+
// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
108+
// CHECK: store ptr null, ptr [[TMP4]], align 8
109+
// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
110+
// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
111+
// CHECK: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0
112+
//
113+
//
114+
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f2v_l20
115+
// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] {
116+
// CHECK: entry:
117+
// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8
118+
// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8
119+
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 1
120+
// CHECK: store i32 6, ptr [[ARRAYIDX]], align 4
121+
// CHECK: ret void
122+
//
123+
//
124+
// CHECK-LABEL: define {{[^@]+}}@_Z2f3v
125+
// CHECK-SAME: () #[[ATTR0]] {
126+
// CHECK: entry:
127+
// CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8
128+
// CHECK: [[TMP1:%.*]] = load ptr, ptr @ptr, align 8
129+
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP1]], i64 0
130+
// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
131+
// CHECK: store ptr @ptr, ptr [[TMP2]], align 8
132+
// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
133+
// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8
134+
// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
135+
// CHECK: store ptr null, ptr [[TMP4]], align 8
136+
// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
137+
// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
138+
// CHECK: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0
139+
//
140+
//
141+
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f3v_l26
142+
// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] {
143+
// CHECK: entry:
144+
// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8
145+
// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8
146+
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 1
147+
// CHECK: store i32 7, ptr [[ARRAYIDX]], align 4
148+
// CHECK: ret void
149+
//
150+
//
151+
// CHECK-LABEL: define {{[^@]+}}@_Z2f4v
152+
// CHECK-SAME: () #[[ATTR0]] {
153+
// CHECK: entry:
154+
// CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8
155+
// CHECK: [[TMP1:%.*]] = load ptr, ptr @ptr, align 8
156+
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 2
157+
// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
158+
// CHECK: store ptr @ptr, ptr [[TMP2]], align 8
159+
// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
160+
// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8
161+
// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
162+
// CHECK: store ptr null, ptr [[TMP4]], align 8
163+
// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
164+
// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
165+
// CHECK: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0
166+
//
167+
//
168+
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f4v_l32
169+
// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] {
170+
// CHECK: entry:
171+
// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8
172+
// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8
173+
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2
174+
// CHECK: store i32 8, ptr [[ARRAYIDX]], align 4
175+
// CHECK: ret void
176+
//
177+
//
178+
// CHECK-LABEL: define {{[^@]+}}@_Z2f5v
179+
// CHECK-SAME: () #[[ATTR0]] {
180+
// CHECK: entry:
181+
// CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8
182+
// CHECK: [[TMP1:%.*]] = load ptr, ptr @ptr, align 8
183+
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 2
184+
// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
185+
// CHECK: store ptr @ptr, ptr [[TMP2]], align 8
186+
// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
187+
// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8
188+
// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
189+
// CHECK: store ptr null, ptr [[TMP4]], align 8
190+
// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
191+
// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
192+
// CHECK: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0
193+
//
194+
//
195+
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f5v_l38
196+
// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] {
197+
// CHECK: entry:
198+
// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8
199+
// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8
200+
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2
201+
// CHECK: store i32 9, ptr [[ARRAYIDX]], align 4
202+
// CHECK: ret void
203+
//
204+
//
205+
// CHECK-LABEL: define {{[^@]+}}@_Z2f6v
206+
// CHECK-SAME: () #[[ATTR0]] {
207+
// CHECK: entry:
208+
// CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8
209+
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2
210+
// CHECK: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
211+
// CHECK: store ptr @ptr, ptr [[TMP1]], align 8
212+
// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
213+
// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP2]], align 8
214+
// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
215+
// CHECK: store ptr null, ptr [[TMP3]], align 8
216+
// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
217+
// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
218+
// CHECK: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 1, ptr [[TMP4]], ptr [[TMP5]], ptr @.offload_sizes.9, ptr @.offload_maptypes.10, ptr null, ptr null)
219+
// CHECK: [[TMP6:%.*]] = load ptr, ptr @ptr, align 8
220+
// CHECK: [[ARRAYIDX1:%.*]] = getelementptr inbounds i32, ptr [[TMP6]], i64 2
221+
// CHECK: store i32 10, ptr [[ARRAYIDX1]], align 4
222+
// CHECK: [[TMP7:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
223+
// CHECK: [[TMP8:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
224+
// CHECK: call void @__tgt_target_data_end_mapper(ptr @[[GLOB1]], i64 -1, i32 1, ptr [[TMP7]], ptr [[TMP8]], ptr @.offload_sizes.9, ptr @.offload_maptypes.10, ptr null, ptr null)
225+
// CHECK: ret void
226+
//
227+
//
228+
// CHECK-LABEL: define {{[^@]+}}@_Z2f7v
229+
// CHECK-SAME: () #[[ATTR0]] {
230+
// CHECK: entry:
231+
// CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8
232+
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2
233+
// CHECK: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
234+
// CHECK: store ptr @ptr, ptr [[TMP1]], align 8
235+
// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
236+
// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP2]], align 8
237+
// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
238+
// CHECK: store ptr null, ptr [[TMP3]], align 8
239+
// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
240+
// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
241+
// CHECK: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1]], i64 -1, i32 1, ptr [[TMP4]], ptr [[TMP5]], ptr @.offload_sizes.11, ptr @.offload_maptypes.12, ptr null, ptr null)
242+
// CHECK: [[TMP6:%.*]] = load ptr, ptr @ptr, align 8
243+
// CHECK: [[ARRAYIDX1:%.*]] = getelementptr inbounds i32, ptr [[TMP6]], i64 2
244+
// CHECK: store i32 11, ptr [[ARRAYIDX1]], align 4
245+
// CHECK: [[TMP7:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
246+
// CHECK: [[TMP8:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
247+
// CHECK: call void @__tgt_target_data_end_mapper(ptr @[[GLOB1]], i64 -1, i32 1, ptr [[TMP7]], ptr [[TMP8]], ptr @.offload_sizes.11, ptr @.offload_maptypes.12, ptr null, ptr null)
248+
// CHECK: ret void
249+
//

0 commit comments

Comments
 (0)