1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
2 // REQUIRES: amdgpu-registered-target
4 // Test without OCML_BASIC_ROUNDED_OPERATIONS
5 // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
6 // RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \
7 // RUN: -internal-isystem %S/Inputs/include \
8 // RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
9 // RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -o - \
10 // RUN: -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,DEFAULT %s
12 // Check that we end up with fast math flags set on intrinsic calls
13 // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
14 // RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \
15 // RUN: -internal-isystem %S/Inputs/include \
16 // RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
17 // RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -ffinite-math-only -o - \
18 // RUN: -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,FINITEONLY %s
21 typedef unsigned long long uint64_t;
23 // CHECK-LABEL: @test___make_mantissa_base8(
25 // CHECK-NEXT: br label [[WHILE_COND_I:%.*]]
26 // CHECK: while.cond.i:
27 // CHECK-NEXT: [[__TAGP_ADDR_0_I:%.*]] = phi ptr [ [[P:%.*]], [[ENTRY:%.*]] ], [ [[__TAGP_ADDR_1_I:%.*]], [[CLEANUP_I:%.*]] ]
28 // CHECK-NEXT: [[__R_0_I:%.*]] = phi i64 [ 0, [[ENTRY]] ], [ [[__R_1_I:%.*]], [[CLEANUP_I]] ]
29 // CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I]], align 1, !tbaa [[TBAA3:![0-9]+]]
30 // CHECK-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP0]], 0
31 // CHECK-NEXT: br i1 [[CMP_NOT_I]], label [[_ZL21__MAKE_MANTISSA_BASE8PKC_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
32 // CHECK: while.body.i:
33 // CHECK-NEXT: [[TMP1:%.*]] = and i8 [[TMP0]], -8
34 // CHECK-NEXT: [[OR_COND_I:%.*]] = icmp eq i8 [[TMP1]], 48
35 // CHECK-NEXT: br i1 [[OR_COND_I]], label [[IF_THEN_I:%.*]], label [[CLEANUP_I]]
37 // CHECK-NEXT: [[MUL_I:%.*]] = shl i64 [[__R_0_I]], 3
38 // CHECK-NEXT: [[CONV5_I:%.*]] = sext i8 [[TMP0]] to i64
39 // CHECK-NEXT: [[ADD_I:%.*]] = add i64 [[MUL_I]], -48
40 // CHECK-NEXT: [[SUB_I:%.*]] = add i64 [[ADD_I]], [[CONV5_I]]
41 // CHECK-NEXT: [[INCDEC_PTR_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I]], i64 1
42 // CHECK-NEXT: br label [[CLEANUP_I]]
44 // CHECK-NEXT: [[__TAGP_ADDR_1_I]] = phi ptr [ [[INCDEC_PTR_I]], [[IF_THEN_I]] ], [ [[__TAGP_ADDR_0_I]], [[WHILE_BODY_I]] ]
45 // CHECK-NEXT: [[__R_1_I]] = phi i64 [ [[SUB_I]], [[IF_THEN_I]] ], [ [[__R_0_I]], [[WHILE_BODY_I]] ]
46 // CHECK-NEXT: br i1 [[OR_COND_I]], label [[WHILE_COND_I]], label [[_ZL21__MAKE_MANTISSA_BASE8PKC_EXIT]], !llvm.loop [[LOOP6:![0-9]+]]
47 // CHECK: _ZL21__make_mantissa_base8PKc.exit:
48 // CHECK-NEXT: [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ]
49 // CHECK-NEXT: ret i64 [[RETVAL_2_I]]
51 extern "C" __device__ uint64_t test___make_mantissa_base8(const char *p) {
52 return __make_mantissa_base8(p);
55 // CHECK-LABEL: @test___make_mantissa_base10(
57 // CHECK-NEXT: br label [[WHILE_COND_I:%.*]]
58 // CHECK: while.cond.i:
59 // CHECK-NEXT: [[__TAGP_ADDR_0_I:%.*]] = phi ptr [ [[P:%.*]], [[ENTRY:%.*]] ], [ [[__TAGP_ADDR_1_I:%.*]], [[CLEANUP_I:%.*]] ]
60 // CHECK-NEXT: [[__R_0_I:%.*]] = phi i64 [ 0, [[ENTRY]] ], [ [[__R_1_I:%.*]], [[CLEANUP_I]] ]
61 // CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I]], align 1, !tbaa [[TBAA3]]
62 // CHECK-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP0]], 0
63 // CHECK-NEXT: br i1 [[CMP_NOT_I]], label [[_ZL22__MAKE_MANTISSA_BASE10PKC_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
64 // CHECK: while.body.i:
65 // CHECK-NEXT: [[TMP1:%.*]] = add i8 [[TMP0]], -48
66 // CHECK-NEXT: [[OR_COND_I:%.*]] = icmp ult i8 [[TMP1]], 10
67 // CHECK-NEXT: br i1 [[OR_COND_I]], label [[IF_THEN_I:%.*]], label [[CLEANUP_I]]
69 // CHECK-NEXT: [[MUL_I:%.*]] = mul i64 [[__R_0_I]], 10
70 // CHECK-NEXT: [[CONV5_I:%.*]] = sext i8 [[TMP0]] to i64
71 // CHECK-NEXT: [[ADD_I:%.*]] = add i64 [[MUL_I]], -48
72 // CHECK-NEXT: [[SUB_I:%.*]] = add i64 [[ADD_I]], [[CONV5_I]]
73 // CHECK-NEXT: [[INCDEC_PTR_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I]], i64 1
74 // CHECK-NEXT: br label [[CLEANUP_I]]
76 // CHECK-NEXT: [[__TAGP_ADDR_1_I]] = phi ptr [ [[INCDEC_PTR_I]], [[IF_THEN_I]] ], [ [[__TAGP_ADDR_0_I]], [[WHILE_BODY_I]] ]
77 // CHECK-NEXT: [[__R_1_I]] = phi i64 [ [[SUB_I]], [[IF_THEN_I]] ], [ [[__R_0_I]], [[WHILE_BODY_I]] ]
78 // CHECK-NEXT: br i1 [[OR_COND_I]], label [[WHILE_COND_I]], label [[_ZL22__MAKE_MANTISSA_BASE10PKC_EXIT]], !llvm.loop [[LOOP9:![0-9]+]]
79 // CHECK: _ZL22__make_mantissa_base10PKc.exit:
80 // CHECK-NEXT: [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ]
81 // CHECK-NEXT: ret i64 [[RETVAL_2_I]]
83 extern "C" __device__ uint64_t test___make_mantissa_base10(const char *p) {
84 return __make_mantissa_base10(p);
87 // CHECK-LABEL: @test___make_mantissa_base16(
89 // CHECK-NEXT: br label [[WHILE_COND_I:%.*]]
90 // CHECK: while.cond.i:
91 // CHECK-NEXT: [[__TAGP_ADDR_0_I:%.*]] = phi ptr [ [[P:%.*]], [[ENTRY:%.*]] ], [ [[__TAGP_ADDR_1_I:%.*]], [[CLEANUP_I:%.*]] ]
92 // CHECK-NEXT: [[__R_0_I:%.*]] = phi i64 [ 0, [[ENTRY]] ], [ [[__R_2_I:%.*]], [[CLEANUP_I]] ]
93 // CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I]], align 1, !tbaa [[TBAA3]]
94 // CHECK-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i8 [[TMP0]], 0
95 // CHECK-NEXT: br i1 [[CMP_NOT_I]], label [[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
96 // CHECK: while.body.i:
97 // CHECK-NEXT: [[TMP1:%.*]] = add i8 [[TMP0]], -48
98 // CHECK-NEXT: [[OR_COND_I:%.*]] = icmp ult i8 [[TMP1]], 10
99 // CHECK-NEXT: br i1 [[OR_COND_I]], label [[IF_END31_I:%.*]], label [[IF_ELSE_I:%.*]]
101 // CHECK-NEXT: [[TMP2:%.*]] = add i8 [[TMP0]], -97
102 // CHECK-NEXT: [[OR_COND33_I:%.*]] = icmp ult i8 [[TMP2]], 6
103 // CHECK-NEXT: br i1 [[OR_COND33_I]], label [[IF_END31_I]], label [[IF_ELSE17_I:%.*]]
104 // CHECK: if.else17.i:
105 // CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP0]], -65
106 // CHECK-NEXT: [[OR_COND34_I:%.*]] = icmp ult i8 [[TMP3]], 6
107 // CHECK-NEXT: br i1 [[OR_COND34_I]], label [[IF_END31_I]], label [[CLEANUP_I]]
108 // CHECK: if.end31.i:
109 // CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I]] ], [ -87, [[IF_ELSE_I]] ], [ -55, [[IF_ELSE17_I]] ]
110 // CHECK-NEXT: [[MUL24_I:%.*]] = shl i64 [[__R_0_I]], 4
111 // CHECK-NEXT: [[CONV25_I:%.*]] = sext i8 [[TMP0]] to i64
112 // CHECK-NEXT: [[ADD26_I:%.*]] = add i64 [[MUL24_I]], [[DOTSINK]]
113 // CHECK-NEXT: [[ADD28_I:%.*]] = add i64 [[ADD26_I]], [[CONV25_I]]
114 // CHECK-NEXT: [[INCDEC_PTR_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I]], i64 1
115 // CHECK-NEXT: br label [[CLEANUP_I]]
117 // CHECK-NEXT: [[__TAGP_ADDR_1_I]] = phi ptr [ [[INCDEC_PTR_I]], [[IF_END31_I]] ], [ [[__TAGP_ADDR_0_I]], [[IF_ELSE17_I]] ]
118 // CHECK-NEXT: [[__R_2_I]] = phi i64 [ [[ADD28_I]], [[IF_END31_I]] ], [ [[__R_0_I]], [[IF_ELSE17_I]] ]
119 // CHECK-NEXT: [[COND_I:%.*]] = phi i1 [ true, [[IF_END31_I]] ], [ false, [[IF_ELSE17_I]] ]
120 // CHECK-NEXT: br i1 [[COND_I]], label [[WHILE_COND_I]], label [[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT]], !llvm.loop [[LOOP10:![0-9]+]]
121 // CHECK: _ZL22__make_mantissa_base16PKc.exit:
122 // CHECK-NEXT: [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ]
123 // CHECK-NEXT: ret i64 [[RETVAL_2_I]]
125 extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
126 return __make_mantissa_base16(p);
129 // CHECK-LABEL: @test___make_mantissa(
130 // CHECK-NEXT: entry:
131 // CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[P:%.*]], align 1, !tbaa [[TBAA3]]
132 // CHECK-NEXT: [[CMP_I:%.*]] = icmp eq i8 [[TMP0]], 48
133 // CHECK-NEXT: br i1 [[CMP_I]], label [[IF_THEN_I:%.*]], label [[WHILE_COND_I33_I:%.*]]
135 // CHECK-NEXT: [[INCDEC_PTR_I:%.*]] = getelementptr inbounds i8, ptr [[P]], i64 1
136 // CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I]], align 1, !tbaa [[TBAA3]]
137 // CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I17_I:%.*]] [
138 // CHECK-NEXT: i8 120, label [[WHILE_COND_I_I_PREHEADER:%.*]]
139 // CHECK-NEXT: i8 88, label [[WHILE_COND_I_I_PREHEADER]]
141 // CHECK: while.cond.i.i.preheader:
142 // CHECK-NEXT: br label [[WHILE_COND_I_I:%.*]]
143 // CHECK: while.cond.i.i:
144 // CHECK-NEXT: [[__TAGP_ADDR_0_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I:%.*]], [[CLEANUP_I_I:%.*]] ], [ [[INCDEC_PTR_I]], [[WHILE_COND_I_I_PREHEADER]] ]
145 // CHECK-NEXT: [[__R_0_I_I:%.*]] = phi i64 [ [[__R_2_I_I:%.*]], [[CLEANUP_I_I]] ], [ 0, [[WHILE_COND_I_I_PREHEADER]] ]
146 // CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I]], align 1, !tbaa [[TBAA3]]
147 // CHECK-NEXT: [[CMP_NOT_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
148 // CHECK-NEXT: br i1 [[CMP_NOT_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT:%.*]], label [[WHILE_BODY_I_I:%.*]]
149 // CHECK: while.body.i.i:
150 // CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
151 // CHECK-NEXT: [[OR_COND_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
152 // CHECK-NEXT: br i1 [[OR_COND_I_I]], label [[IF_END31_I_I:%.*]], label [[IF_ELSE_I_I:%.*]]
153 // CHECK: if.else.i.i:
154 // CHECK-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
155 // CHECK-NEXT: [[OR_COND33_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
156 // CHECK-NEXT: br i1 [[OR_COND33_I_I]], label [[IF_END31_I_I]], label [[IF_ELSE17_I_I:%.*]]
157 // CHECK: if.else17.i.i:
158 // CHECK-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
159 // CHECK-NEXT: [[OR_COND34_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
160 // CHECK-NEXT: br i1 [[OR_COND34_I_I]], label [[IF_END31_I_I]], label [[CLEANUP_I_I]]
161 // CHECK: if.end31.i.i:
162 // CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I]] ], [ -87, [[IF_ELSE_I_I]] ], [ -55, [[IF_ELSE17_I_I]] ]
163 // CHECK-NEXT: [[MUL24_I_I:%.*]] = shl i64 [[__R_0_I_I]], 4
164 // CHECK-NEXT: [[CONV25_I_I:%.*]] = sext i8 [[TMP2]] to i64
165 // CHECK-NEXT: [[ADD26_I_I:%.*]] = add i64 [[MUL24_I_I]], [[DOTSINK]]
166 // CHECK-NEXT: [[ADD28_I_I:%.*]] = add i64 [[ADD26_I_I]], [[CONV25_I_I]]
167 // CHECK-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I_I]], i64 1
168 // CHECK-NEXT: br label [[CLEANUP_I_I]]
169 // CHECK: cleanup.i.i:
170 // CHECK-NEXT: [[__TAGP_ADDR_1_I_I]] = phi ptr [ [[INCDEC_PTR_I_I]], [[IF_END31_I_I]] ], [ [[__TAGP_ADDR_0_I_I]], [[IF_ELSE17_I_I]] ]
171 // CHECK-NEXT: [[__R_2_I_I]] = phi i64 [ [[ADD28_I_I]], [[IF_END31_I_I]] ], [ [[__R_0_I_I]], [[IF_ELSE17_I_I]] ]
172 // CHECK-NEXT: [[COND_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I]] ], [ false, [[IF_ELSE17_I_I]] ]
173 // CHECK-NEXT: br i1 [[COND_I_I]], label [[WHILE_COND_I_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP10]]
174 // CHECK: while.cond.i17.i:
175 // CHECK-NEXT: [[__TAGP_ADDR_0_I14_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I:%.*]], [[CLEANUP_I28_I:%.*]] ], [ [[INCDEC_PTR_I]], [[IF_THEN_I]] ]
176 // CHECK-NEXT: [[__R_0_I15_I:%.*]] = phi i64 [ [[__R_1_I27_I:%.*]], [[CLEANUP_I28_I]] ], [ 0, [[IF_THEN_I]] ]
177 // CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I]], align 1, !tbaa [[TBAA3]]
178 // CHECK-NEXT: [[CMP_NOT_I16_I:%.*]] = icmp eq i8 [[TMP6]], 0
179 // CHECK-NEXT: br i1 [[CMP_NOT_I16_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I19_I:%.*]]
180 // CHECK: while.body.i19.i:
181 // CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
182 // CHECK-NEXT: [[OR_COND_I18_I:%.*]] = icmp eq i8 [[TMP7]], 48
183 // CHECK-NEXT: br i1 [[OR_COND_I18_I]], label [[IF_THEN_I25_I:%.*]], label [[CLEANUP_I28_I]]
184 // CHECK: if.then.i25.i:
185 // CHECK-NEXT: [[MUL_I20_I:%.*]] = shl i64 [[__R_0_I15_I]], 3
186 // CHECK-NEXT: [[CONV5_I21_I:%.*]] = sext i8 [[TMP6]] to i64
187 // CHECK-NEXT: [[ADD_I22_I:%.*]] = add i64 [[MUL_I20_I]], -48
188 // CHECK-NEXT: [[SUB_I23_I:%.*]] = add i64 [[ADD_I22_I]], [[CONV5_I21_I]]
189 // CHECK-NEXT: [[INCDEC_PTR_I24_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I14_I]], i64 1
190 // CHECK-NEXT: br label [[CLEANUP_I28_I]]
191 // CHECK: cleanup.i28.i:
192 // CHECK-NEXT: [[__TAGP_ADDR_1_I26_I]] = phi ptr [ [[INCDEC_PTR_I24_I]], [[IF_THEN_I25_I]] ], [ [[__TAGP_ADDR_0_I14_I]], [[WHILE_BODY_I19_I]] ]
193 // CHECK-NEXT: [[__R_1_I27_I]] = phi i64 [ [[SUB_I23_I]], [[IF_THEN_I25_I]] ], [ [[__R_0_I15_I]], [[WHILE_BODY_I19_I]] ]
194 // CHECK-NEXT: br i1 [[OR_COND_I18_I]], label [[WHILE_COND_I17_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP6]]
195 // CHECK: while.cond.i33.i:
196 // CHECK-NEXT: [[__TAGP_ADDR_0_I30_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I42_I:%.*]], [[CLEANUP_I44_I:%.*]] ], [ [[P]], [[ENTRY:%.*]] ]
197 // CHECK-NEXT: [[__R_0_I31_I:%.*]] = phi i64 [ [[__R_1_I43_I:%.*]], [[CLEANUP_I44_I]] ], [ 0, [[ENTRY]] ]
198 // CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I]], align 1, !tbaa [[TBAA3]]
199 // CHECK-NEXT: [[CMP_NOT_I32_I:%.*]] = icmp eq i8 [[TMP8]], 0
200 // CHECK-NEXT: br i1 [[CMP_NOT_I32_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], label [[WHILE_BODY_I35_I:%.*]]
201 // CHECK: while.body.i35.i:
202 // CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
203 // CHECK-NEXT: [[OR_COND_I34_I:%.*]] = icmp ult i8 [[TMP9]], 10
204 // CHECK-NEXT: br i1 [[OR_COND_I34_I]], label [[IF_THEN_I41_I:%.*]], label [[CLEANUP_I44_I]]
205 // CHECK: if.then.i41.i:
206 // CHECK-NEXT: [[MUL_I36_I:%.*]] = mul i64 [[__R_0_I31_I]], 10
207 // CHECK-NEXT: [[CONV5_I37_I:%.*]] = sext i8 [[TMP8]] to i64
208 // CHECK-NEXT: [[ADD_I38_I:%.*]] = add i64 [[MUL_I36_I]], -48
209 // CHECK-NEXT: [[SUB_I39_I:%.*]] = add i64 [[ADD_I38_I]], [[CONV5_I37_I]]
210 // CHECK-NEXT: [[INCDEC_PTR_I40_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I]], i64 1
211 // CHECK-NEXT: br label [[CLEANUP_I44_I]]
212 // CHECK: cleanup.i44.i:
213 // CHECK-NEXT: [[__TAGP_ADDR_1_I42_I]] = phi ptr [ [[INCDEC_PTR_I40_I]], [[IF_THEN_I41_I]] ], [ [[__TAGP_ADDR_0_I30_I]], [[WHILE_BODY_I35_I]] ]
214 // CHECK-NEXT: [[__R_1_I43_I]] = phi i64 [ [[SUB_I39_I]], [[IF_THEN_I41_I]] ], [ [[__R_0_I31_I]], [[WHILE_BODY_I35_I]] ]
215 // CHECK-NEXT: br i1 [[OR_COND_I34_I]], label [[WHILE_COND_I33_I]], label [[_ZL15__MAKE_MANTISSAPKC_EXIT]], !llvm.loop [[LOOP9]]
216 // CHECK: _ZL15__make_mantissaPKc.exit:
217 // CHECK-NEXT: [[RETVAL_0_I:%.*]] = phi i64 [ 0, [[CLEANUP_I28_I]] ], [ [[__R_0_I15_I]], [[WHILE_COND_I17_I]] ], [ 0, [[CLEANUP_I_I]] ], [ [[__R_0_I_I]], [[WHILE_COND_I_I]] ], [ 0, [[CLEANUP_I44_I]] ], [ [[__R_0_I31_I]], [[WHILE_COND_I33_I]] ]
218 // CHECK-NEXT: ret i64 [[RETVAL_0_I]]
220 extern "C" __device__ uint64_t test___make_mantissa(const char *p) {
221 return __make_mantissa(p);
224 // CHECK-LABEL: @test_abs(
225 // CHECK-NEXT: entry:
226 // CHECK-NEXT: [[SUB_I:%.*]] = tail call i32 @llvm.abs.i32(i32 [[X:%.*]], i1 true)
227 // CHECK-NEXT: ret i32 [[SUB_I]]
229 extern "C" __device__ int test_abs(int x) {
233 // CHECK-LABEL: @test_labs(
234 // CHECK-NEXT: entry:
235 // CHECK-NEXT: [[SUB_I:%.*]] = tail call i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true)
236 // CHECK-NEXT: ret i64 [[SUB_I]]
238 extern "C" __device__ long test_labs(long x) {
242 // CHECK-LABEL: @test_llabs(
243 // CHECK-NEXT: entry:
244 // CHECK-NEXT: [[SUB_I:%.*]] = tail call i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true)
245 // CHECK-NEXT: ret i64 [[SUB_I]]
247 extern "C" __device__ long long test_llabs(long x) {
251 // DEFAULT-LABEL: @test_acosf(
252 // DEFAULT-NEXT: entry:
253 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_acos_f32(float noundef [[X:%.*]]) #[[ATTR13:[0-9]+]]
254 // DEFAULT-NEXT: ret float [[CALL_I]]
256 // FINITEONLY-LABEL: @test_acosf(
257 // FINITEONLY-NEXT: entry:
258 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_acos_f32(float noundef [[X:%.*]]) #[[ATTR13:[0-9]+]]
259 // FINITEONLY-NEXT: ret float [[CALL_I]]
261 extern "C" __device__ float test_acosf(float x) {
265 // DEFAULT-LABEL: @test_acos(
266 // DEFAULT-NEXT: entry:
267 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_acos_f64(double noundef [[X:%.*]]) #[[ATTR13]]
268 // DEFAULT-NEXT: ret double [[CALL_I]]
270 // FINITEONLY-LABEL: @test_acos(
271 // FINITEONLY-NEXT: entry:
272 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_acos_f64(double noundef [[X:%.*]]) #[[ATTR13]]
273 // FINITEONLY-NEXT: ret double [[CALL_I]]
275 extern "C" __device__ double test_acos(double x) {
279 // DEFAULT-LABEL: @test_acoshf(
280 // DEFAULT-NEXT: entry:
281 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_acosh_f32(float noundef [[X:%.*]]) #[[ATTR14:[0-9]+]]
282 // DEFAULT-NEXT: ret float [[CALL_I]]
284 // FINITEONLY-LABEL: @test_acoshf(
285 // FINITEONLY-NEXT: entry:
286 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_acosh_f32(float noundef [[X:%.*]]) #[[ATTR14:[0-9]+]]
287 // FINITEONLY-NEXT: ret float [[CALL_I]]
289 extern "C" __device__ float test_acoshf(float x) {
293 // DEFAULT-LABEL: @test_acosh(
294 // DEFAULT-NEXT: entry:
295 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_acosh_f64(double noundef [[X:%.*]]) #[[ATTR14]]
296 // DEFAULT-NEXT: ret double [[CALL_I]]
298 // FINITEONLY-LABEL: @test_acosh(
299 // FINITEONLY-NEXT: entry:
300 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_acosh_f64(double noundef [[X:%.*]]) #[[ATTR14]]
301 // FINITEONLY-NEXT: ret double [[CALL_I]]
303 extern "C" __device__ double test_acosh(double x) {
307 // DEFAULT-LABEL: @test_asinf(
308 // DEFAULT-NEXT: entry:
309 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_asin_f32(float noundef [[X:%.*]]) #[[ATTR13]]
310 // DEFAULT-NEXT: ret float [[CALL_I]]
312 // FINITEONLY-LABEL: @test_asinf(
313 // FINITEONLY-NEXT: entry:
314 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_asin_f32(float noundef [[X:%.*]]) #[[ATTR13]]
315 // FINITEONLY-NEXT: ret float [[CALL_I]]
317 extern "C" __device__ float test_asinf(float x) {
321 // DEFAULT-LABEL: @test_asin(
322 // DEFAULT-NEXT: entry:
323 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_asin_f64(double noundef [[X:%.*]]) #[[ATTR13]]
324 // DEFAULT-NEXT: ret double [[CALL_I]]
326 // FINITEONLY-LABEL: @test_asin(
327 // FINITEONLY-NEXT: entry:
328 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_asin_f64(double noundef [[X:%.*]]) #[[ATTR13]]
329 // FINITEONLY-NEXT: ret double [[CALL_I]]
331 extern "C" __device__ double test_asin(double x) {
336 // DEFAULT-LABEL: @test_asinhf(
337 // DEFAULT-NEXT: entry:
338 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_asinh_f32(float noundef [[X:%.*]]) #[[ATTR14]]
339 // DEFAULT-NEXT: ret float [[CALL_I]]
341 // FINITEONLY-LABEL: @test_asinhf(
342 // FINITEONLY-NEXT: entry:
343 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_asinh_f32(float noundef [[X:%.*]]) #[[ATTR14]]
344 // FINITEONLY-NEXT: ret float [[CALL_I]]
346 extern "C" __device__ float test_asinhf(float x) {
350 // DEFAULT-LABEL: @test_asinh(
351 // DEFAULT-NEXT: entry:
352 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_asinh_f64(double noundef [[X:%.*]]) #[[ATTR14]]
353 // DEFAULT-NEXT: ret double [[CALL_I]]
355 // FINITEONLY-LABEL: @test_asinh(
356 // FINITEONLY-NEXT: entry:
357 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_asinh_f64(double noundef [[X:%.*]]) #[[ATTR14]]
358 // FINITEONLY-NEXT: ret double [[CALL_I]]
360 extern "C" __device__ double test_asinh(double x) {
364 // DEFAULT-LABEL: @test_atan2f(
365 // DEFAULT-NEXT: entry:
366 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_atan2_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
367 // DEFAULT-NEXT: ret float [[CALL_I]]
369 // FINITEONLY-LABEL: @test_atan2f(
370 // FINITEONLY-NEXT: entry:
371 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_atan2_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
372 // FINITEONLY-NEXT: ret float [[CALL_I]]
374 extern "C" __device__ float test_atan2f(float x, float y) {
378 // DEFAULT-LABEL: @test_atan2(
379 // DEFAULT-NEXT: entry:
380 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_atan2_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
381 // DEFAULT-NEXT: ret double [[CALL_I]]
383 // FINITEONLY-LABEL: @test_atan2(
384 // FINITEONLY-NEXT: entry:
385 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_atan2_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
386 // FINITEONLY-NEXT: ret double [[CALL_I]]
388 extern "C" __device__ double test_atan2(double x, double y) {
392 // DEFAULT-LABEL: @test_atanf(
393 // DEFAULT-NEXT: entry:
394 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_atan_f32(float noundef [[X:%.*]]) #[[ATTR13]]
395 // DEFAULT-NEXT: ret float [[CALL_I]]
397 // FINITEONLY-LABEL: @test_atanf(
398 // FINITEONLY-NEXT: entry:
399 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_atan_f32(float noundef [[X:%.*]]) #[[ATTR13]]
400 // FINITEONLY-NEXT: ret float [[CALL_I]]
402 extern "C" __device__ float test_atanf(float x) {
406 // DEFAULT-LABEL: @test_atan(
407 // DEFAULT-NEXT: entry:
408 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_atan_f64(double noundef [[X:%.*]]) #[[ATTR13]]
409 // DEFAULT-NEXT: ret double [[CALL_I]]
411 // FINITEONLY-LABEL: @test_atan(
412 // FINITEONLY-NEXT: entry:
413 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_atan_f64(double noundef [[X:%.*]]) #[[ATTR13]]
414 // FINITEONLY-NEXT: ret double [[CALL_I]]
416 extern "C" __device__ double test_atan(double x) {
420 // DEFAULT-LABEL: @test_atanhf(
421 // DEFAULT-NEXT: entry:
422 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_atanh_f32(float noundef [[X:%.*]]) #[[ATTR14]]
423 // DEFAULT-NEXT: ret float [[CALL_I]]
425 // FINITEONLY-LABEL: @test_atanhf(
426 // FINITEONLY-NEXT: entry:
427 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_atanh_f32(float noundef [[X:%.*]]) #[[ATTR14]]
428 // FINITEONLY-NEXT: ret float [[CALL_I]]
430 extern "C" __device__ float test_atanhf(float x) {
434 // DEFAULT-LABEL: @test_atanh(
435 // DEFAULT-NEXT: entry:
436 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_atanh_f64(double noundef [[X:%.*]]) #[[ATTR14]]
437 // DEFAULT-NEXT: ret double [[CALL_I]]
439 // FINITEONLY-LABEL: @test_atanh(
440 // FINITEONLY-NEXT: entry:
441 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_atanh_f64(double noundef [[X:%.*]]) #[[ATTR14]]
442 // FINITEONLY-NEXT: ret double [[CALL_I]]
444 extern "C" __device__ double test_atanh(double x) {
448 // DEFAULT-LABEL: @test_cbrtf(
449 // DEFAULT-NEXT: entry:
450 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cbrt_f32(float noundef [[X:%.*]]) #[[ATTR14]]
451 // DEFAULT-NEXT: ret float [[CALL_I]]
453 // FINITEONLY-LABEL: @test_cbrtf(
454 // FINITEONLY-NEXT: entry:
455 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cbrt_f32(float noundef [[X:%.*]]) #[[ATTR14]]
456 // FINITEONLY-NEXT: ret float [[CALL_I]]
458 extern "C" __device__ float test_cbrtf(float x) {
462 // DEFAULT-LABEL: @test_cbrt(
463 // DEFAULT-NEXT: entry:
464 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cbrt_f64(double noundef [[X:%.*]]) #[[ATTR14]]
465 // DEFAULT-NEXT: ret double [[CALL_I]]
467 // FINITEONLY-LABEL: @test_cbrt(
468 // FINITEONLY-NEXT: entry:
469 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cbrt_f64(double noundef [[X:%.*]]) #[[ATTR14]]
470 // FINITEONLY-NEXT: ret double [[CALL_I]]
472 extern "C" __device__ double test_cbrt(double x) {
476 // DEFAULT-LABEL: @test_ceilf(
477 // DEFAULT-NEXT: entry:
478 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ceil_f32(float noundef [[X:%.*]]) #[[ATTR13]]
479 // DEFAULT-NEXT: ret float [[CALL_I]]
481 // FINITEONLY-LABEL: @test_ceilf(
482 // FINITEONLY-NEXT: entry:
483 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ceil_f32(float noundef [[X:%.*]]) #[[ATTR13]]
484 // FINITEONLY-NEXT: ret float [[CALL_I]]
486 extern "C" __device__ float test_ceilf(float x) {
490 // DEFAULT-LABEL: @test_ceil(
491 // DEFAULT-NEXT: entry:
492 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ceil_f64(double noundef [[X:%.*]]) #[[ATTR13]]
493 // DEFAULT-NEXT: ret double [[CALL_I]]
495 // FINITEONLY-LABEL: @test_ceil(
496 // FINITEONLY-NEXT: entry:
497 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ceil_f64(double noundef [[X:%.*]]) #[[ATTR13]]
498 // FINITEONLY-NEXT: ret double [[CALL_I]]
500 extern "C" __device__ double test_ceil(double x) {
504 // DEFAULT-LABEL: @test_copysignf(
505 // DEFAULT-NEXT: entry:
506 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_copysign_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
507 // DEFAULT-NEXT: ret float [[CALL_I]]
509 // FINITEONLY-LABEL: @test_copysignf(
510 // FINITEONLY-NEXT: entry:
511 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_copysign_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
512 // FINITEONLY-NEXT: ret float [[CALL_I]]
514 extern "C" __device__ float test_copysignf(float x, float y) {
515 return copysignf(x, y);
518 // DEFAULT-LABEL: @test_copysign(
519 // DEFAULT-NEXT: entry:
520 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_copysign_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
521 // DEFAULT-NEXT: ret double [[CALL_I]]
523 // FINITEONLY-LABEL: @test_copysign(
524 // FINITEONLY-NEXT: entry:
525 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_copysign_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
526 // FINITEONLY-NEXT: ret double [[CALL_I]]
528 extern "C" __device__ double test_copysign(double x, double y) {
529 return copysign(x, y);
532 // DEFAULT-LABEL: @test_cosf(
533 // DEFAULT-NEXT: entry:
534 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR15:[0-9]+]]
535 // DEFAULT-NEXT: ret float [[CALL_I]]
537 // FINITEONLY-LABEL: @test_cosf(
538 // FINITEONLY-NEXT: entry:
539 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR15:[0-9]+]]
540 // FINITEONLY-NEXT: ret float [[CALL_I]]
542 extern "C" __device__ float test_cosf(float x) {
546 // DEFAULT-LABEL: @test_cos(
547 // DEFAULT-NEXT: entry:
548 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cos_f64(double noundef [[X:%.*]]) #[[ATTR15]]
549 // DEFAULT-NEXT: ret double [[CALL_I]]
551 // FINITEONLY-LABEL: @test_cos(
552 // FINITEONLY-NEXT: entry:
553 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cos_f64(double noundef [[X:%.*]]) #[[ATTR15]]
554 // FINITEONLY-NEXT: ret double [[CALL_I]]
556 extern "C" __device__ double test_cos(double x) {
560 // DEFAULT-LABEL: @test_coshf(
561 // DEFAULT-NEXT: entry:
562 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cosh_f32(float noundef [[X:%.*]]) #[[ATTR14]]
563 // DEFAULT-NEXT: ret float [[CALL_I]]
565 // FINITEONLY-LABEL: @test_coshf(
566 // FINITEONLY-NEXT: entry:
567 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cosh_f32(float noundef [[X:%.*]]) #[[ATTR14]]
568 // FINITEONLY-NEXT: ret float [[CALL_I]]
570 extern "C" __device__ float test_coshf(float x) {
574 // DEFAULT-LABEL: @test_cosh(
575 // DEFAULT-NEXT: entry:
576 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cosh_f64(double noundef [[X:%.*]]) #[[ATTR14]]
577 // DEFAULT-NEXT: ret double [[CALL_I]]
579 // FINITEONLY-LABEL: @test_cosh(
580 // FINITEONLY-NEXT: entry:
581 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cosh_f64(double noundef [[X:%.*]]) #[[ATTR14]]
582 // FINITEONLY-NEXT: ret double [[CALL_I]]
584 extern "C" __device__ double test_cosh(double x) {
588 // DEFAULT-LABEL: @test_cospif(
589 // DEFAULT-NEXT: entry:
590 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_cospi_f32(float noundef [[X:%.*]]) #[[ATTR15]]
591 // DEFAULT-NEXT: ret float [[CALL_I]]
593 // FINITEONLY-LABEL: @test_cospif(
594 // FINITEONLY-NEXT: entry:
595 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cospi_f32(float noundef [[X:%.*]]) #[[ATTR15]]
596 // FINITEONLY-NEXT: ret float [[CALL_I]]
598 extern "C" __device__ float test_cospif(float x) {
602 // DEFAULT-LABEL: @test_cospi(
603 // DEFAULT-NEXT: entry:
604 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_cospi_f64(double noundef [[X:%.*]]) #[[ATTR15]]
605 // DEFAULT-NEXT: ret double [[CALL_I]]
607 // FINITEONLY-LABEL: @test_cospi(
608 // FINITEONLY-NEXT: entry:
609 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cospi_f64(double noundef [[X:%.*]]) #[[ATTR15]]
610 // FINITEONLY-NEXT: ret double [[CALL_I]]
612 extern "C" __device__ double test_cospi(double x) {
616 // DEFAULT-LABEL: @test_cyl_bessel_i0f(
617 // DEFAULT-NEXT: entry:
618 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_i0_f32(float noundef [[X:%.*]]) #[[ATTR15]]
619 // DEFAULT-NEXT: ret float [[CALL_I]]
621 // FINITEONLY-LABEL: @test_cyl_bessel_i0f(
622 // FINITEONLY-NEXT: entry:
623 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_i0_f32(float noundef [[X:%.*]]) #[[ATTR15]]
624 // FINITEONLY-NEXT: ret float [[CALL_I]]
626 extern "C" __device__ float test_cyl_bessel_i0f(float x) {
627 return cyl_bessel_i0f(x);
630 // DEFAULT-LABEL: @test_cyl_bessel_i0(
631 // DEFAULT-NEXT: entry:
632 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_i0_f64(double noundef [[X:%.*]]) #[[ATTR15]]
633 // DEFAULT-NEXT: ret double [[CALL_I]]
635 // FINITEONLY-LABEL: @test_cyl_bessel_i0(
636 // FINITEONLY-NEXT: entry:
637 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_i0_f64(double noundef [[X:%.*]]) #[[ATTR15]]
638 // FINITEONLY-NEXT: ret double [[CALL_I]]
640 extern "C" __device__ double test_cyl_bessel_i0(double x) {
641 return cyl_bessel_i0(x);
644 // DEFAULT-LABEL: @test_cyl_bessel_i1f(
645 // DEFAULT-NEXT: entry:
646 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_i1_f32(float noundef [[X:%.*]]) #[[ATTR15]]
647 // DEFAULT-NEXT: ret float [[CALL_I]]
649 // FINITEONLY-LABEL: @test_cyl_bessel_i1f(
650 // FINITEONLY-NEXT: entry:
651 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_i1_f32(float noundef [[X:%.*]]) #[[ATTR15]]
652 // FINITEONLY-NEXT: ret float [[CALL_I]]
654 extern "C" __device__ float test_cyl_bessel_i1f(float x) {
655 return cyl_bessel_i1f(x);
658 // DEFAULT-LABEL: @test_cyl_bessel_i1(
659 // DEFAULT-NEXT: entry:
660 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_i1_f64(double noundef [[X:%.*]]) #[[ATTR15]]
661 // DEFAULT-NEXT: ret double [[CALL_I]]
663 // FINITEONLY-LABEL: @test_cyl_bessel_i1(
664 // FINITEONLY-NEXT: entry:
665 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_i1_f64(double noundef [[X:%.*]]) #[[ATTR15]]
666 // FINITEONLY-NEXT: ret double [[CALL_I]]
668 extern "C" __device__ double test_cyl_bessel_i1(double x) {
669 return cyl_bessel_i1(x);
672 // DEFAULT-LABEL: @test_erfcf(
673 // DEFAULT-NEXT: entry:
674 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_erfc_f32(float noundef [[X:%.*]]) #[[ATTR14]]
675 // DEFAULT-NEXT: ret float [[CALL_I]]
677 // FINITEONLY-LABEL: @test_erfcf(
678 // FINITEONLY-NEXT: entry:
679 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_erfc_f32(float noundef [[X:%.*]]) #[[ATTR14]]
680 // FINITEONLY-NEXT: ret float [[CALL_I]]
682 extern "C" __device__ float test_erfcf(float x) {
686 // DEFAULT-LABEL: @test_erfc(
687 // DEFAULT-NEXT: entry:
688 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_erfc_f64(double noundef [[X:%.*]]) #[[ATTR14]]
689 // DEFAULT-NEXT: ret double [[CALL_I]]
691 // FINITEONLY-LABEL: @test_erfc(
692 // FINITEONLY-NEXT: entry:
693 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_erfc_f64(double noundef [[X:%.*]]) #[[ATTR14]]
694 // FINITEONLY-NEXT: ret double [[CALL_I]]
696 extern "C" __device__ double test_erfc(double x) {
700 // DEFAULT-LABEL: @test_erfinvf(
701 // DEFAULT-NEXT: entry:
702 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_erfinv_f32(float noundef [[X:%.*]]) #[[ATTR14]]
703 // DEFAULT-NEXT: ret float [[CALL_I]]
705 // FINITEONLY-LABEL: @test_erfinvf(
706 // FINITEONLY-NEXT: entry:
707 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_erfinv_f32(float noundef [[X:%.*]]) #[[ATTR14]]
708 // FINITEONLY-NEXT: ret float [[CALL_I]]
710 extern "C" __device__ float test_erfinvf(float x) {
714 // DEFAULT-LABEL: @test_erfinv(
715 // DEFAULT-NEXT: entry:
716 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_erfinv_f64(double noundef [[X:%.*]]) #[[ATTR14]]
717 // DEFAULT-NEXT: ret double [[CALL_I]]
719 // FINITEONLY-LABEL: @test_erfinv(
720 // FINITEONLY-NEXT: entry:
721 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_erfinv_f64(double noundef [[X:%.*]]) #[[ATTR14]]
722 // FINITEONLY-NEXT: ret double [[CALL_I]]
724 extern "C" __device__ double test_erfinv(double x) {
728 // DEFAULT-LABEL: @test_exp10f(
729 // DEFAULT-NEXT: entry:
730 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_exp10_f32(float noundef [[X:%.*]]) #[[ATTR14]]
731 // DEFAULT-NEXT: ret float [[CALL_I]]
733 // FINITEONLY-LABEL: @test_exp10f(
734 // FINITEONLY-NEXT: entry:
735 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_exp10_f32(float noundef [[X:%.*]]) #[[ATTR14]]
736 // FINITEONLY-NEXT: ret float [[CALL_I]]
738 extern "C" __device__ float test_exp10f(float x) {
742 // DEFAULT-LABEL: @test_exp10(
743 // DEFAULT-NEXT: entry:
744 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp10_f64(double noundef [[X:%.*]]) #[[ATTR14]]
745 // DEFAULT-NEXT: ret double [[CALL_I]]
747 // FINITEONLY-LABEL: @test_exp10(
748 // FINITEONLY-NEXT: entry:
749 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_exp10_f64(double noundef [[X:%.*]]) #[[ATTR14]]
750 // FINITEONLY-NEXT: ret double [[CALL_I]]
752 extern "C" __device__ double test_exp10(double x) {
756 // DEFAULT-LABEL: @test_exp2f(
757 // DEFAULT-NEXT: entry:
758 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_exp2_f32(float noundef [[X:%.*]]) #[[ATTR14]]
759 // DEFAULT-NEXT: ret float [[CALL_I]]
761 // FINITEONLY-LABEL: @test_exp2f(
762 // FINITEONLY-NEXT: entry:
763 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_exp2_f32(float noundef [[X:%.*]]) #[[ATTR14]]
764 // FINITEONLY-NEXT: ret float [[CALL_I]]
766 extern "C" __device__ float test_exp2f(float x) {
771 // DEFAULT-LABEL: @test_exp2(
772 // DEFAULT-NEXT: entry:
773 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR14]]
774 // DEFAULT-NEXT: ret double [[CALL_I]]
776 // FINITEONLY-LABEL: @test_exp2(
777 // FINITEONLY-NEXT: entry:
778 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR14]]
779 // FINITEONLY-NEXT: ret double [[CALL_I]]
781 extern "C" __device__ double test_exp2(double x) {
785 // DEFAULT-LABEL: @test_expf(
786 // DEFAULT-NEXT: entry:
787 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_exp_f32(float noundef [[X:%.*]]) #[[ATTR14]]
788 // DEFAULT-NEXT: ret float [[CALL_I]]
790 // FINITEONLY-LABEL: @test_expf(
791 // FINITEONLY-NEXT: entry:
792 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_exp_f32(float noundef [[X:%.*]]) #[[ATTR14]]
793 // FINITEONLY-NEXT: ret float [[CALL_I]]
795 extern "C" __device__ float test_expf(float x) {
799 // DEFAULT-LABEL: @test_exp(
800 // DEFAULT-NEXT: entry:
801 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_exp_f64(double noundef [[X:%.*]]) #[[ATTR14]]
802 // DEFAULT-NEXT: ret double [[CALL_I]]
804 // FINITEONLY-LABEL: @test_exp(
805 // FINITEONLY-NEXT: entry:
806 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_exp_f64(double noundef [[X:%.*]]) #[[ATTR14]]
807 // FINITEONLY-NEXT: ret double [[CALL_I]]
809 extern "C" __device__ double test_exp(double x) {
813 // DEFAULT-LABEL: @test_expm1f(
814 // DEFAULT-NEXT: entry:
815 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_expm1_f32(float noundef [[X:%.*]]) #[[ATTR14]]
816 // DEFAULT-NEXT: ret float [[CALL_I]]
818 // FINITEONLY-LABEL: @test_expm1f(
819 // FINITEONLY-NEXT: entry:
820 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_expm1_f32(float noundef [[X:%.*]]) #[[ATTR14]]
821 // FINITEONLY-NEXT: ret float [[CALL_I]]
823 extern "C" __device__ float test_expm1f(float x) {
827 // DEFAULT-LABEL: @test_expm1(
828 // DEFAULT-NEXT: entry:
829 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_expm1_f64(double noundef [[X:%.*]]) #[[ATTR14]]
830 // DEFAULT-NEXT: ret double [[CALL_I]]
832 // FINITEONLY-LABEL: @test_expm1(
833 // FINITEONLY-NEXT: entry:
834 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_expm1_f64(double noundef [[X:%.*]]) #[[ATTR14]]
835 // FINITEONLY-NEXT: ret double [[CALL_I]]
837 extern "C" __device__ double test_expm1(double x) {
841 // DEFAULT-LABEL: @test_fabsf(
842 // DEFAULT-NEXT: entry:
843 // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.fabs.f32(float [[X:%.*]])
844 // DEFAULT-NEXT: ret float [[TMP0]]
846 // FINITEONLY-LABEL: @test_fabsf(
847 // FINITEONLY-NEXT: entry:
848 // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.fabs.f32(float [[X:%.*]])
849 // FINITEONLY-NEXT: ret float [[TMP0]]
851 extern "C" __device__ float test_fabsf(float x) {
855 // DEFAULT-LABEL: @test_fabs(
856 // DEFAULT-NEXT: entry:
857 // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.fabs.f64(double [[X:%.*]])
858 // DEFAULT-NEXT: ret double [[TMP0]]
860 // FINITEONLY-LABEL: @test_fabs(
861 // FINITEONLY-NEXT: entry:
862 // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.fabs.f64(double [[X:%.*]])
863 // FINITEONLY-NEXT: ret double [[TMP0]]
865 extern "C" __device__ double test_fabs(double x) {
869 // DEFAULT-LABEL: @test_fdimf(
870 // DEFAULT-NEXT: entry:
871 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fdim_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
872 // DEFAULT-NEXT: ret float [[CALL_I]]
874 // FINITEONLY-LABEL: @test_fdimf(
875 // FINITEONLY-NEXT: entry:
876 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fdim_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
877 // FINITEONLY-NEXT: ret float [[CALL_I]]
879 extern "C" __device__ float test_fdimf(float x, float y) {
883 // DEFAULT-LABEL: @test_fdim(
884 // DEFAULT-NEXT: entry:
885 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fdim_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
886 // DEFAULT-NEXT: ret double [[CALL_I]]
888 // FINITEONLY-LABEL: @test_fdim(
889 // FINITEONLY-NEXT: entry:
890 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fdim_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
891 // FINITEONLY-NEXT: ret double [[CALL_I]]
893 extern "C" __device__ double test_fdim(double x, double y) {
897 // DEFAULT-LABEL: @test_fdividef(
898 // DEFAULT-NEXT: entry:
899 // DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float [[X:%.*]], [[Y:%.*]]
900 // DEFAULT-NEXT: ret float [[DIV_I]]
902 // FINITEONLY-LABEL: @test_fdividef(
903 // FINITEONLY-NEXT: entry:
904 // FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract float [[X:%.*]], [[Y:%.*]]
905 // FINITEONLY-NEXT: ret float [[DIV_I]]
907 extern "C" __device__ float test_fdividef(float x, float y) {
908 return fdividef(x, y);
911 // DEFAULT-LABEL: @test_floorf(
912 // DEFAULT-NEXT: entry:
913 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_floor_f32(float noundef [[X:%.*]]) #[[ATTR13]]
914 // DEFAULT-NEXT: ret float [[CALL_I]]
916 // FINITEONLY-LABEL: @test_floorf(
917 // FINITEONLY-NEXT: entry:
918 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_floor_f32(float noundef [[X:%.*]]) #[[ATTR13]]
919 // FINITEONLY-NEXT: ret float [[CALL_I]]
921 extern "C" __device__ float test_floorf(float x) {
925 // DEFAULT-LABEL: @test_floor(
926 // DEFAULT-NEXT: entry:
927 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_floor_f64(double noundef [[X:%.*]]) #[[ATTR13]]
928 // DEFAULT-NEXT: ret double [[CALL_I]]
930 // FINITEONLY-LABEL: @test_floor(
931 // FINITEONLY-NEXT: entry:
932 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_floor_f64(double noundef [[X:%.*]]) #[[ATTR13]]
933 // FINITEONLY-NEXT: ret double [[CALL_I]]
935 extern "C" __device__ double test_floor(double x) {
939 // DEFAULT-LABEL: @test_fmaf(
940 // DEFAULT-NEXT: entry:
941 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR13]]
942 // DEFAULT-NEXT: ret float [[CALL_I]]
944 // FINITEONLY-LABEL: @test_fmaf(
945 // FINITEONLY-NEXT: entry:
946 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR13]]
947 // FINITEONLY-NEXT: ret float [[CALL_I]]
949 extern "C" __device__ float test_fmaf(float x, float y, float z) {
950 return fmaf(x, y, z);
953 // DEFAULT-LABEL: @test_fma(
954 // DEFAULT-NEXT: entry:
955 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR13]]
956 // DEFAULT-NEXT: ret double [[CALL_I]]
958 // FINITEONLY-LABEL: @test_fma(
959 // FINITEONLY-NEXT: entry:
960 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR13]]
961 // FINITEONLY-NEXT: ret double [[CALL_I]]
963 extern "C" __device__ double test_fma(double x, double y, double z) {
967 // DEFAULT-LABEL: @test_fma_rn(
968 // DEFAULT-NEXT: entry:
969 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR13]]
970 // DEFAULT-NEXT: ret double [[CALL_I]]
972 // FINITEONLY-LABEL: @test_fma_rn(
973 // FINITEONLY-NEXT: entry:
974 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR13]]
975 // FINITEONLY-NEXT: ret double [[CALL_I]]
977 extern "C" __device__ double test_fma_rn(double x, double y, double z) {
978 return __fma_rn(x, y, z);
981 // DEFAULT-LABEL: @test_fmaxf(
982 // DEFAULT-NEXT: entry:
983 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
984 // DEFAULT-NEXT: ret float [[CALL_I]]
986 // FINITEONLY-LABEL: @test_fmaxf(
987 // FINITEONLY-NEXT: entry:
988 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
989 // FINITEONLY-NEXT: ret float [[CALL_I]]
991 extern "C" __device__ float test_fmaxf(float x, float y) {
995 // DEFAULT-LABEL: @test_fmax(
996 // DEFAULT-NEXT: entry:
997 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
998 // DEFAULT-NEXT: ret double [[CALL_I]]
1000 // FINITEONLY-LABEL: @test_fmax(
1001 // FINITEONLY-NEXT: entry:
1002 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
1003 // FINITEONLY-NEXT: ret double [[CALL_I]]
1005 extern "C" __device__ double test_fmax(double x, double y) {
1009 // DEFAULT-LABEL: @test_fminf(
1010 // DEFAULT-NEXT: entry:
1011 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
1012 // DEFAULT-NEXT: ret float [[CALL_I]]
1014 // FINITEONLY-LABEL: @test_fminf(
1015 // FINITEONLY-NEXT: entry:
1016 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
1017 // FINITEONLY-NEXT: ret float [[CALL_I]]
1019 extern "C" __device__ float test_fminf(float x, float y) {
1023 // DEFAULT-LABEL: @test_fmin(
1024 // DEFAULT-NEXT: entry:
1025 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
1026 // DEFAULT-NEXT: ret double [[CALL_I]]
1028 // FINITEONLY-LABEL: @test_fmin(
1029 // FINITEONLY-NEXT: entry:
1030 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
1031 // FINITEONLY-NEXT: ret double [[CALL_I]]
1033 extern "C" __device__ double test_fmin(double x, double y) {
1037 // DEFAULT-LABEL: @test_fmodf(
1038 // DEFAULT-NEXT: entry:
1039 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fmod_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
1040 // DEFAULT-NEXT: ret float [[CALL_I]]
1042 // FINITEONLY-LABEL: @test_fmodf(
1043 // FINITEONLY-NEXT: entry:
1044 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fmod_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
1045 // FINITEONLY-NEXT: ret float [[CALL_I]]
1047 extern "C" __device__ float test_fmodf(float x, float y) {
1051 // DEFAULT-LABEL: @test_fmod(
1052 // DEFAULT-NEXT: entry:
1053 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fmod_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
1054 // DEFAULT-NEXT: ret double [[CALL_I]]
1056 // FINITEONLY-LABEL: @test_fmod(
1057 // FINITEONLY-NEXT: entry:
1058 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fmod_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
1059 // FINITEONLY-NEXT: ret double [[CALL_I]]
1061 extern "C" __device__ double test_fmod(double x, double y) {
1065 // DEFAULT-LABEL: @test_frexpf(
1066 // DEFAULT-NEXT: entry:
1067 // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
1068 // DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16:[0-9]+]]
1069 // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_frexp_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
1070 // DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11:![0-9]+]]
1071 // DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11]]
1072 // DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
1073 // DEFAULT-NEXT: ret float [[CALL_I]]
1075 // FINITEONLY-LABEL: @test_frexpf(
1076 // FINITEONLY-NEXT: entry:
1077 // FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
1078 // FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16:[0-9]+]]
1079 // FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_frexp_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
1080 // FINITEONLY-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11:![0-9]+]]
1081 // FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11]]
1082 // FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
1083 // FINITEONLY-NEXT: ret float [[CALL_I]]
1085 extern "C" __device__ float test_frexpf(float x, int* y) {
1086 return frexpf(x, y);
1089 // DEFAULT-LABEL: @test_frexp(
1090 // DEFAULT-NEXT: entry:
1091 // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
1092 // DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
1093 // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_frexp_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
1094 // DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]]
1095 // DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11]]
1096 // DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
1097 // DEFAULT-NEXT: ret double [[CALL_I]]
1099 // FINITEONLY-LABEL: @test_frexp(
1100 // FINITEONLY-NEXT: entry:
1101 // FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
1102 // FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
1103 // FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_frexp_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
1104 // FINITEONLY-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]]
1105 // FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA11]]
1106 // FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
1107 // FINITEONLY-NEXT: ret double [[CALL_I]]
1109 extern "C" __device__ double test_frexp(double x, int* y) {
1113 // DEFAULT-LABEL: @test_hypotf(
1114 // DEFAULT-NEXT: entry:
1115 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_hypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
1116 // DEFAULT-NEXT: ret float [[CALL_I]]
1118 // FINITEONLY-LABEL: @test_hypotf(
1119 // FINITEONLY-NEXT: entry:
1120 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_hypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
1121 // FINITEONLY-NEXT: ret float [[CALL_I]]
1123 extern "C" __device__ float test_hypotf(float x, float y) {
1124 return hypotf(x, y);
1127 // DEFAULT-LABEL: @test_hypot(
1128 // DEFAULT-NEXT: entry:
1129 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_hypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
1130 // DEFAULT-NEXT: ret double [[CALL_I]]
1132 // FINITEONLY-LABEL: @test_hypot(
1133 // FINITEONLY-NEXT: entry:
1134 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_hypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
1135 // FINITEONLY-NEXT: ret double [[CALL_I]]
1137 extern "C" __device__ double test_hypot(double x, double y) {
1141 // CHECK-LABEL: @test_ilogbf(
1142 // CHECK-NEXT: entry:
1143 // CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f32(float noundef [[X:%.*]]) #[[ATTR13:[0-9]+]]
1144 // CHECK-NEXT: ret i32 [[CALL_I]]
1146 extern "C" __device__ int test_ilogbf(float x) {
1150 // CHECK-LABEL: @test_ilogb(
1151 // CHECK-NEXT: entry:
1152 // CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f64(double noundef [[X:%.*]]) #[[ATTR13]]
1153 // CHECK-NEXT: ret i32 [[CALL_I]]
1155 extern "C" __device__ int test_ilogb(double x) {
1159 // CHECK-LABEL: @test___finitef(
1160 // CHECK-NEXT: entry:
1161 // CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isfinite_f32(float noundef [[X:%.*]]) #[[ATTR13]]
1162 // CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
1163 // CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
1164 // CHECK-NEXT: ret i32 [[CONV]]
1166 extern "C" __device__ BOOL_TYPE test___finitef(float x) {
1167 return __finitef(x);
1170 // CHECK-LABEL: @test___finite(
1171 // CHECK-NEXT: entry:
1172 // CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isfinite_f64(double noundef [[X:%.*]]) #[[ATTR13]]
1173 // CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
1174 // CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
1175 // CHECK-NEXT: ret i32 [[CONV]]
1177 extern "C" __device__ BOOL_TYPE test___finite(double x) {
1181 // CHECK-LABEL: @test___isinff(
1182 // CHECK-NEXT: entry:
1183 // CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f32(float noundef [[X:%.*]]) #[[ATTR13]]
1184 // CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
1185 // CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
1186 // CHECK-NEXT: ret i32 [[CONV]]
1188 extern "C" __device__ BOOL_TYPE test___isinff(float x) {
1192 // CHECK-LABEL: @test___isinf(
1193 // CHECK-NEXT: entry:
1194 // CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f64(double noundef [[X:%.*]]) #[[ATTR13]]
1195 // CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
1196 // CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
1197 // CHECK-NEXT: ret i32 [[CONV]]
1199 extern "C" __device__ BOOL_TYPE test___isinf(double x) {
1203 // CHECK-LABEL: @test___isnanf(
1204 // CHECK-NEXT: entry:
1205 // CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f32(float noundef [[X:%.*]]) #[[ATTR13]]
1206 // CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
1207 // CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
1208 // CHECK-NEXT: ret i32 [[CONV]]
1210 extern "C" __device__ BOOL_TYPE test___isnanf(float x) {
1214 // CHECK-LABEL: @test___isnan(
1215 // CHECK-NEXT: entry:
1216 // CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f64(double noundef [[X:%.*]]) #[[ATTR13]]
1217 // CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
1218 // CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
1219 // CHECK-NEXT: ret i32 [[CONV]]
1221 extern "C" __device__ BOOL_TYPE test___isnan(double x) {
1225 // DEFAULT-LABEL: @test_j0f(
1226 // DEFAULT-NEXT: entry:
1227 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[X:%.*]]) #[[ATTR15]]
1228 // DEFAULT-NEXT: ret float [[CALL_I]]
1230 // FINITEONLY-LABEL: @test_j0f(
1231 // FINITEONLY-NEXT: entry:
1232 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[X:%.*]]) #[[ATTR15]]
1233 // FINITEONLY-NEXT: ret float [[CALL_I]]
1235 extern "C" __device__ float test_j0f(float x) {
1239 // DEFAULT-LABEL: @test_j0(
1240 // DEFAULT-NEXT: entry:
1241 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[X:%.*]]) #[[ATTR15]]
1242 // DEFAULT-NEXT: ret double [[CALL_I]]
1244 // FINITEONLY-LABEL: @test_j0(
1245 // FINITEONLY-NEXT: entry:
1246 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[X:%.*]]) #[[ATTR15]]
1247 // FINITEONLY-NEXT: ret double [[CALL_I]]
1249 extern "C" __device__ double test_j0(double x) {
1253 // DEFAULT-LABEL: @test_j1f(
1254 // DEFAULT-NEXT: entry:
1255 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[X:%.*]]) #[[ATTR15]]
1256 // DEFAULT-NEXT: ret float [[CALL_I]]
1258 // FINITEONLY-LABEL: @test_j1f(
1259 // FINITEONLY-NEXT: entry:
1260 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[X:%.*]]) #[[ATTR15]]
1261 // FINITEONLY-NEXT: ret float [[CALL_I]]
1263 extern "C" __device__ float test_j1f(float x) {
1267 // DEFAULT-LABEL: @test_j1(
1268 // DEFAULT-NEXT: entry:
1269 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[X:%.*]]) #[[ATTR15]]
1270 // DEFAULT-NEXT: ret double [[CALL_I]]
1272 // FINITEONLY-LABEL: @test_j1(
1273 // FINITEONLY-NEXT: entry:
1274 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[X:%.*]]) #[[ATTR15]]
1275 // FINITEONLY-NEXT: ret double [[CALL_I]]
1277 extern "C" __device__ double test_j1(double x) {
1281 // DEFAULT-LABEL: @test_jnf(
1282 // DEFAULT-NEXT: entry:
1283 // DEFAULT-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
1284 // DEFAULT-NEXT: i32 0, label [[IF_THEN_I:%.*]]
1285 // DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
1287 // DEFAULT: if.then.i:
1288 // DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
1289 // DEFAULT-NEXT: br label [[_ZL3JNFIF_EXIT:%.*]]
1290 // DEFAULT: if.then2.i:
1291 // DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
1292 // DEFAULT-NEXT: br label [[_ZL3JNFIF_EXIT]]
1293 // DEFAULT: if.end4.i:
1294 // DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR15]]
1295 // DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
1296 // DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
1297 // DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
1298 // DEFAULT: for.body.i:
1299 // DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
1300 // DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
1301 // DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
1302 // DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
1303 // DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
1304 // DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
1305 // DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]]
1306 // DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]]
1307 // DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
1308 // DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
1309 // DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP13:![0-9]+]]
1310 // DEFAULT: _ZL3jnfif.exit:
1311 // DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
1312 // DEFAULT-NEXT: ret float [[RETVAL_0_I]]
1314 // FINITEONLY-LABEL: @test_jnf(
1315 // FINITEONLY-NEXT: entry:
1316 // FINITEONLY-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
1317 // FINITEONLY-NEXT: i32 0, label [[IF_THEN_I:%.*]]
1318 // FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
1319 // FINITEONLY-NEXT: ]
1320 // FINITEONLY: if.then.i:
1321 // FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
1322 // FINITEONLY-NEXT: br label [[_ZL3JNFIF_EXIT:%.*]]
1323 // FINITEONLY: if.then2.i:
1324 // FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
1325 // FINITEONLY-NEXT: br label [[_ZL3JNFIF_EXIT]]
1326 // FINITEONLY: if.end4.i:
1327 // FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR15]]
1328 // FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR15]]
1329 // FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
1330 // FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
1331 // FINITEONLY: for.body.i:
1332 // FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
1333 // FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
1334 // FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
1335 // FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
1336 // FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
1337 // FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract float [[CONV_I]], [[Y]]
1338 // FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_025_I]], [[DIV_I]]
1339 // FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_024_I]]
1340 // FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
1341 // FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
1342 // FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP13:![0-9]+]]
1343 // FINITEONLY: _ZL3jnfif.exit:
1344 // FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
1345 // FINITEONLY-NEXT: ret float [[RETVAL_0_I]]
1347 extern "C" __device__ float test_jnf(int x, float y) {
1351 // DEFAULT-LABEL: @test_jn(
1352 // DEFAULT-NEXT: entry:
1353 // DEFAULT-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
1354 // DEFAULT-NEXT: i32 0, label [[IF_THEN_I:%.*]]
1355 // DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
1357 // DEFAULT: if.then.i:
1358 // DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
1359 // DEFAULT-NEXT: br label [[_ZL2JNID_EXIT:%.*]]
1360 // DEFAULT: if.then2.i:
1361 // DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
1362 // DEFAULT-NEXT: br label [[_ZL2JNID_EXIT]]
1363 // DEFAULT: if.end4.i:
1364 // DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR15]]
1365 // DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
1366 // DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
1367 // DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
1368 // DEFAULT: for.body.i:
1369 // DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
1370 // DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
1371 // DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
1372 // DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
1373 // DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
1374 // DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
1375 // DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]]
1376 // DEFAULT-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]]
1377 // DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
1378 // DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
1379 // DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]]
1380 // DEFAULT: _ZL2jnid.exit:
1381 // DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
1382 // DEFAULT-NEXT: ret double [[RETVAL_0_I]]
1384 // FINITEONLY-LABEL: @test_jn(
1385 // FINITEONLY-NEXT: entry:
1386 // FINITEONLY-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
1387 // FINITEONLY-NEXT: i32 0, label [[IF_THEN_I:%.*]]
1388 // FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
1389 // FINITEONLY-NEXT: ]
1390 // FINITEONLY: if.then.i:
1391 // FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
1392 // FINITEONLY-NEXT: br label [[_ZL2JNID_EXIT:%.*]]
1393 // FINITEONLY: if.then2.i:
1394 // FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
1395 // FINITEONLY-NEXT: br label [[_ZL2JNID_EXIT]]
1396 // FINITEONLY: if.end4.i:
1397 // FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR15]]
1398 // FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR15]]
1399 // FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
1400 // FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
1401 // FINITEONLY: for.body.i:
1402 // FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
1403 // FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
1404 // FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
1405 // FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
1406 // FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
1407 // FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract double [[CONV_I]], [[Y]]
1408 // FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_025_I]], [[DIV_I]]
1409 // FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_024_I]]
1410 // FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
1411 // FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
1412 // FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]]
1413 // FINITEONLY: _ZL2jnid.exit:
1414 // FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
1415 // FINITEONLY-NEXT: ret double [[RETVAL_0_I]]
1417 extern "C" __device__ double test_jn(int x, double y) {
1421 // DEFAULT-LABEL: @test_ldexpf(
1422 // DEFAULT-NEXT: entry:
1423 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ldexp_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]]
1424 // DEFAULT-NEXT: ret float [[CALL_I]]
1426 // FINITEONLY-LABEL: @test_ldexpf(
1427 // FINITEONLY-NEXT: entry:
1428 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ldexp_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]]
1429 // FINITEONLY-NEXT: ret float [[CALL_I]]
1431 extern "C" __device__ float test_ldexpf(float x, int y) {
1432 return ldexpf(x, y);
1435 // DEFAULT-LABEL: @test_ldexp(
1436 // DEFAULT-NEXT: entry:
1437 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ldexp_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]]
1438 // DEFAULT-NEXT: ret double [[CALL_I]]
1440 // FINITEONLY-LABEL: @test_ldexp(
1441 // FINITEONLY-NEXT: entry:
1442 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ldexp_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]]
1443 // FINITEONLY-NEXT: ret double [[CALL_I]]
1445 extern "C" __device__ double test_ldexp(double x, int y) {
1449 // DEFAULT-LABEL: @test_lgammaf(
1450 // DEFAULT-NEXT: entry:
1451 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_lgamma_f32(float noundef [[X:%.*]]) #[[ATTR15]]
1452 // DEFAULT-NEXT: ret float [[CALL_I]]
1454 // FINITEONLY-LABEL: @test_lgammaf(
1455 // FINITEONLY-NEXT: entry:
1456 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_lgamma_f32(float noundef [[X:%.*]]) #[[ATTR15]]
1457 // FINITEONLY-NEXT: ret float [[CALL_I]]
1459 extern "C" __device__ float test_lgammaf(float x) {
1463 // DEFAULT-LABEL: @test_lgamma(
1464 // DEFAULT-NEXT: entry:
1465 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_lgamma_f64(double noundef [[X:%.*]]) #[[ATTR15]]
1466 // DEFAULT-NEXT: ret double [[CALL_I]]
1468 // FINITEONLY-LABEL: @test_lgamma(
1469 // FINITEONLY-NEXT: entry:
1470 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_lgamma_f64(double noundef [[X:%.*]]) #[[ATTR15]]
1471 // FINITEONLY-NEXT: ret double [[CALL_I]]
1473 extern "C" __device__ double test_lgamma(double x) {
1477 // DEFAULT-LABEL: @test_llrintf(
1478 // DEFAULT-NEXT: entry:
1479 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR13]]
1480 // DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
1481 // DEFAULT-NEXT: ret i64 [[CONV_I]]
1483 // FINITEONLY-LABEL: @test_llrintf(
1484 // FINITEONLY-NEXT: entry:
1485 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR13]]
1486 // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
1487 // FINITEONLY-NEXT: ret i64 [[CONV_I]]
1489 extern "C" __device__ long long int test_llrintf(float x) {
1493 // DEFAULT-LABEL: @test_llrint(
1494 // DEFAULT-NEXT: entry:
1495 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR13]]
1496 // DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
1497 // DEFAULT-NEXT: ret i64 [[CONV_I]]
1499 // FINITEONLY-LABEL: @test_llrint(
1500 // FINITEONLY-NEXT: entry:
1501 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR13]]
1502 // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
1503 // FINITEONLY-NEXT: ret i64 [[CONV_I]]
1505 extern "C" __device__ long long int test_llrint(double x) {
1509 // DEFAULT-LABEL: @test_llroundf(
1510 // DEFAULT-NEXT: entry:
1511 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR13]]
1512 // DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
1513 // DEFAULT-NEXT: ret i64 [[CONV_I]]
1515 // FINITEONLY-LABEL: @test_llroundf(
1516 // FINITEONLY-NEXT: entry:
1517 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR13]]
1518 // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
1519 // FINITEONLY-NEXT: ret i64 [[CONV_I]]
1521 extern "C" __device__ long long int test_llroundf(float x) {
1525 // DEFAULT-LABEL: @test_llround(
1526 // DEFAULT-NEXT: entry:
1527 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR13]]
1528 // DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
1529 // DEFAULT-NEXT: ret i64 [[CONV_I]]
1531 // FINITEONLY-LABEL: @test_llround(
1532 // FINITEONLY-NEXT: entry:
1533 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR13]]
1534 // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
1535 // FINITEONLY-NEXT: ret i64 [[CONV_I]]
1537 extern "C" __device__ long long int test_llround(double x) {
1541 // DEFAULT-LABEL: @test_log10f(
1542 // DEFAULT-NEXT: entry:
1543 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_log10_f32(float noundef [[X:%.*]]) #[[ATTR14]]
1544 // DEFAULT-NEXT: ret float [[CALL_I]]
1546 // FINITEONLY-LABEL: @test_log10f(
1547 // FINITEONLY-NEXT: entry:
1548 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_log10_f32(float noundef [[X:%.*]]) #[[ATTR14]]
1549 // FINITEONLY-NEXT: ret float [[CALL_I]]
1551 extern "C" __device__ float test_log10f(float x) {
1555 // DEFAULT-LABEL: @test_log10(
1556 // DEFAULT-NEXT: entry:
1557 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_log10_f64(double noundef [[X:%.*]]) #[[ATTR14]]
1558 // DEFAULT-NEXT: ret double [[CALL_I]]
1560 // FINITEONLY-LABEL: @test_log10(
1561 // FINITEONLY-NEXT: entry:
1562 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_log10_f64(double noundef [[X:%.*]]) #[[ATTR14]]
1563 // FINITEONLY-NEXT: ret double [[CALL_I]]
1565 extern "C" __device__ double test_log10(double x) {
1569 // DEFAULT-LABEL: @test_log1pf(
1570 // DEFAULT-NEXT: entry:
1571 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_log1p_f32(float noundef [[X:%.*]]) #[[ATTR14]]
1572 // DEFAULT-NEXT: ret float [[CALL_I]]
1574 // FINITEONLY-LABEL: @test_log1pf(
1575 // FINITEONLY-NEXT: entry:
1576 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_log1p_f32(float noundef [[X:%.*]]) #[[ATTR14]]
1577 // FINITEONLY-NEXT: ret float [[CALL_I]]
1579 extern "C" __device__ float test_log1pf(float x) {
1583 // DEFAULT-LABEL: @test_log1p(
1584 // DEFAULT-NEXT: entry:
1585 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_log1p_f64(double noundef [[X:%.*]]) #[[ATTR14]]
1586 // DEFAULT-NEXT: ret double [[CALL_I]]
1588 // FINITEONLY-LABEL: @test_log1p(
1589 // FINITEONLY-NEXT: entry:
1590 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_log1p_f64(double noundef [[X:%.*]]) #[[ATTR14]]
1591 // FINITEONLY-NEXT: ret double [[CALL_I]]
1593 extern "C" __device__ double test_log1p(double x) {
1597 // DEFAULT-LABEL: @test_log2f(
1598 // DEFAULT-NEXT: entry:
1599 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_log2_f32(float noundef [[X:%.*]]) #[[ATTR14]]
1600 // DEFAULT-NEXT: ret float [[CALL_I]]
1602 // FINITEONLY-LABEL: @test_log2f(
1603 // FINITEONLY-NEXT: entry:
1604 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_log2_f32(float noundef [[X:%.*]]) #[[ATTR14]]
1605 // FINITEONLY-NEXT: ret float [[CALL_I]]
1607 extern "C" __device__ float test_log2f(float x) {
1611 // DEFAULT-LABEL: @test_log2(
1612 // DEFAULT-NEXT: entry:
1613 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_log2_f64(double noundef [[X:%.*]]) #[[ATTR14]]
1614 // DEFAULT-NEXT: ret double [[CALL_I]]
1616 // FINITEONLY-LABEL: @test_log2(
1617 // FINITEONLY-NEXT: entry:
1618 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_log2_f64(double noundef [[X:%.*]]) #[[ATTR14]]
1619 // FINITEONLY-NEXT: ret double [[CALL_I]]
1621 extern "C" __device__ double test_log2(double x) {
1625 // DEFAULT-LABEL: @test_logbf(
1626 // DEFAULT-NEXT: entry:
1627 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_logb_f32(float noundef [[X:%.*]]) #[[ATTR13]]
1628 // DEFAULT-NEXT: ret float [[CALL_I]]
1630 // FINITEONLY-LABEL: @test_logbf(
1631 // FINITEONLY-NEXT: entry:
1632 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_logb_f32(float noundef [[X:%.*]]) #[[ATTR13]]
1633 // FINITEONLY-NEXT: ret float [[CALL_I]]
1635 extern "C" __device__ float test_logbf(float x) {
1639 // DEFAULT-LABEL: @test_logb(
1640 // DEFAULT-NEXT: entry:
1641 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_logb_f64(double noundef [[X:%.*]]) #[[ATTR13]]
1642 // DEFAULT-NEXT: ret double [[CALL_I]]
1644 // FINITEONLY-LABEL: @test_logb(
1645 // FINITEONLY-NEXT: entry:
1646 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_logb_f64(double noundef [[X:%.*]]) #[[ATTR13]]
1647 // FINITEONLY-NEXT: ret double [[CALL_I]]
1649 extern "C" __device__ double test_logb(double x) {
1653 // DEFAULT-LABEL: @test_lrintf(
1654 // DEFAULT-NEXT: entry:
1655 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR13]]
1656 // DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
1657 // DEFAULT-NEXT: ret i64 [[CONV_I]]
1659 // FINITEONLY-LABEL: @test_lrintf(
1660 // FINITEONLY-NEXT: entry:
1661 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR13]]
1662 // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
1663 // FINITEONLY-NEXT: ret i64 [[CONV_I]]
1665 extern "C" __device__ long int test_lrintf(float x) {
1669 // DEFAULT-LABEL: @test_lrint(
1670 // DEFAULT-NEXT: entry:
1671 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR13]]
1672 // DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
1673 // DEFAULT-NEXT: ret i64 [[CONV_I]]
1675 // FINITEONLY-LABEL: @test_lrint(
1676 // FINITEONLY-NEXT: entry:
1677 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR13]]
1678 // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
1679 // FINITEONLY-NEXT: ret i64 [[CONV_I]]
1681 extern "C" __device__ long int test_lrint(double x) {
1685 // DEFAULT-LABEL: @test_lroundf(
1686 // DEFAULT-NEXT: entry:
1687 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR13]]
1688 // DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
1689 // DEFAULT-NEXT: ret i64 [[CONV_I]]
1691 // FINITEONLY-LABEL: @test_lroundf(
1692 // FINITEONLY-NEXT: entry:
1693 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR13]]
1694 // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
1695 // FINITEONLY-NEXT: ret i64 [[CONV_I]]
1697 extern "C" __device__ long int test_lroundf(float x) {
1701 // DEFAULT-LABEL: @test_lround(
1702 // DEFAULT-NEXT: entry:
1703 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR13]]
1704 // DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
1705 // DEFAULT-NEXT: ret i64 [[CONV_I]]
1707 // FINITEONLY-LABEL: @test_lround(
1708 // FINITEONLY-NEXT: entry:
1709 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR13]]
1710 // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
1711 // FINITEONLY-NEXT: ret i64 [[CONV_I]]
1713 extern "C" __device__ long int test_lround(double x) {
1717 // DEFAULT-LABEL: @test_modff(
1718 // DEFAULT-NEXT: entry:
1719 // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
1720 // DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
1721 // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_modf_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
1722 // DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15:![0-9]+]]
1723 // DEFAULT-NEXT: store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]]
1724 // DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
1725 // DEFAULT-NEXT: ret float [[CALL_I]]
1727 // FINITEONLY-LABEL: @test_modff(
1728 // FINITEONLY-NEXT: entry:
1729 // FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
1730 // FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
1731 // FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_modf_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
1732 // FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15:![0-9]+]]
1733 // FINITEONLY-NEXT: store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]]
1734 // FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
1735 // FINITEONLY-NEXT: ret float [[CALL_I]]
1737 extern "C" __device__ float test_modff(float x, float* y) {
1741 // DEFAULT-LABEL: @test_modf(
1742 // DEFAULT-NEXT: entry:
1743 // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
1744 // DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
1745 // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
1746 // DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17:![0-9]+]]
1747 // DEFAULT-NEXT: store double [[TMP0]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]]
1748 // DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
1749 // DEFAULT-NEXT: ret double [[CALL_I]]
1751 // FINITEONLY-LABEL: @test_modf(
1752 // FINITEONLY-NEXT: entry:
1753 // FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
1754 // FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
1755 // FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
1756 // FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17:![0-9]+]]
1757 // FINITEONLY-NEXT: store double [[TMP0]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]]
1758 // FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
1759 // FINITEONLY-NEXT: ret double [[CALL_I]]
1761 extern "C" __device__ double test_modf(double x, double* y) {
1765 // CHECK-LABEL: @test_nanf(
1766 // CHECK-NEXT: entry:
1767 // CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa [[TBAA3]]
1768 // CHECK-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
1769 // CHECK-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I33_I_I:%.*]]
1770 // CHECK: if.then.i.i:
1771 // CHECK-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr [[TAG]], i64 1
1772 // CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[TBAA3]]
1773 // CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I17_I_I:%.*]] [
1774 // CHECK-NEXT: i8 120, label [[WHILE_COND_I_I_I_PREHEADER:%.*]]
1775 // CHECK-NEXT: i8 88, label [[WHILE_COND_I_I_I_PREHEADER]]
1777 // CHECK: while.cond.i.i.i.preheader:
1778 // CHECK-NEXT: br label [[WHILE_COND_I_I_I:%.*]]
1779 // CHECK: while.cond.i.i.i:
1780 // CHECK-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I_I_I_PREHEADER]] ]
1781 // CHECK-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[WHILE_COND_I_I_I_PREHEADER]] ]
1782 // CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA3]]
1783 // CHECK-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
1784 // CHECK-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL4NANFPKC_EXIT:%.*]], label [[WHILE_BODY_I_I_I:%.*]]
1785 // CHECK: while.body.i.i.i:
1786 // CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
1787 // CHECK-NEXT: [[OR_COND_I_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
1788 // CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]]
1789 // CHECK: if.else.i.i.i:
1790 // CHECK-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
1791 // CHECK-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
1792 // CHECK-NEXT: br i1 [[OR_COND33_I_I_I]], label [[IF_END31_I_I_I]], label [[IF_ELSE17_I_I_I:%.*]]
1793 // CHECK: if.else17.i.i.i:
1794 // CHECK-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
1795 // CHECK-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
1796 // CHECK-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label [[CLEANUP_I_I_I]]
1797 // CHECK: if.end31.i.i.i:
1798 // CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
1799 // CHECK-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 4
1800 // CHECK-NEXT: [[CONV25_I_I_I:%.*]] = sext i8 [[TMP2]] to i64
1801 // CHECK-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
1802 // CHECK-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
1803 // CHECK-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I_I_I]], i64 1
1804 // CHECK-NEXT: br label [[CLEANUP_I_I_I]]
1805 // CHECK: cleanup.i.i.i:
1806 // CHECK-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[IF_ELSE17_I_I_I]] ]
1807 // CHECK-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__R_0_I_I_I]], [[IF_ELSE17_I_I_I]] ]
1808 // CHECK-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [ false, [[IF_ELSE17_I_I_I]] ]
1809 // CHECK-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP10]]
1810 // CHECK: while.cond.i17.i.i:
1811 // CHECK-NEXT: [[__TAGP_ADDR_0_I14_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I_I:%.*]], [[CLEANUP_I28_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ]
1812 // CHECK-NEXT: [[__R_0_I15_I_I:%.*]] = phi i64 [ [[__R_1_I27_I_I:%.*]], [[CLEANUP_I28_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
1813 // CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I_I]], align 1, !tbaa [[TBAA3]]
1814 // CHECK-NEXT: [[CMP_NOT_I16_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
1815 // CHECK-NEXT: br i1 [[CMP_NOT_I16_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I19_I_I:%.*]]
1816 // CHECK: while.body.i19.i.i:
1817 // CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
1818 // CHECK-NEXT: [[OR_COND_I18_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
1819 // CHECK-NEXT: br i1 [[OR_COND_I18_I_I]], label [[IF_THEN_I25_I_I:%.*]], label [[CLEANUP_I28_I_I]]
1820 // CHECK: if.then.i25.i.i:
1821 // CHECK-NEXT: [[MUL_I20_I_I:%.*]] = shl i64 [[__R_0_I15_I_I]], 3
1822 // CHECK-NEXT: [[CONV5_I21_I_I:%.*]] = sext i8 [[TMP6]] to i64
1823 // CHECK-NEXT: [[ADD_I22_I_I:%.*]] = add i64 [[MUL_I20_I_I]], -48
1824 // CHECK-NEXT: [[SUB_I23_I_I:%.*]] = add i64 [[ADD_I22_I_I]], [[CONV5_I21_I_I]]
1825 // CHECK-NEXT: [[INCDEC_PTR_I24_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I14_I_I]], i64 1
1826 // CHECK-NEXT: br label [[CLEANUP_I28_I_I]]
1827 // CHECK: cleanup.i28.i.i:
1828 // CHECK-NEXT: [[__TAGP_ADDR_1_I26_I_I]] = phi ptr [ [[INCDEC_PTR_I24_I_I]], [[IF_THEN_I25_I_I]] ], [ [[__TAGP_ADDR_0_I14_I_I]], [[WHILE_BODY_I19_I_I]] ]
1829 // CHECK-NEXT: [[__R_1_I27_I_I]] = phi i64 [ [[SUB_I23_I_I]], [[IF_THEN_I25_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_BODY_I19_I_I]] ]
1830 // CHECK-NEXT: br i1 [[OR_COND_I18_I_I]], label [[WHILE_COND_I17_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP6]]
1831 // CHECK: while.cond.i33.i.i:
1832 // CHECK-NEXT: [[__TAGP_ADDR_0_I30_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I42_I_I:%.*]], [[CLEANUP_I44_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ]
1833 // CHECK-NEXT: [[__R_0_I31_I_I:%.*]] = phi i64 [ [[__R_1_I43_I_I:%.*]], [[CLEANUP_I44_I_I]] ], [ 0, [[ENTRY]] ]
1834 // CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I_I]], align 1, !tbaa [[TBAA3]]
1835 // CHECK-NEXT: [[CMP_NOT_I32_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
1836 // CHECK-NEXT: br i1 [[CMP_NOT_I32_I_I]], label [[_ZL4NANFPKC_EXIT]], label [[WHILE_BODY_I35_I_I:%.*]]
1837 // CHECK: while.body.i35.i.i:
1838 // CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
1839 // CHECK-NEXT: [[OR_COND_I34_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
1840 // CHECK-NEXT: br i1 [[OR_COND_I34_I_I]], label [[IF_THEN_I41_I_I:%.*]], label [[CLEANUP_I44_I_I]]
1841 // CHECK: if.then.i41.i.i:
1842 // CHECK-NEXT: [[MUL_I36_I_I:%.*]] = mul i64 [[__R_0_I31_I_I]], 10
1843 // CHECK-NEXT: [[CONV5_I37_I_I:%.*]] = sext i8 [[TMP8]] to i64
1844 // CHECK-NEXT: [[ADD_I38_I_I:%.*]] = add i64 [[MUL_I36_I_I]], -48
1845 // CHECK-NEXT: [[SUB_I39_I_I:%.*]] = add i64 [[ADD_I38_I_I]], [[CONV5_I37_I_I]]
1846 // CHECK-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I_I]], i64 1
1847 // CHECK-NEXT: br label [[CLEANUP_I44_I_I]]
1848 // CHECK: cleanup.i44.i.i:
1849 // CHECK-NEXT: [[__TAGP_ADDR_1_I42_I_I]] = phi ptr [ [[INCDEC_PTR_I40_I_I]], [[IF_THEN_I41_I_I]] ], [ [[__TAGP_ADDR_0_I30_I_I]], [[WHILE_BODY_I35_I_I]] ]
1850 // CHECK-NEXT: [[__R_1_I43_I_I]] = phi i64 [ [[SUB_I39_I_I]], [[IF_THEN_I41_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_BODY_I35_I_I]] ]
1851 // CHECK-NEXT: br i1 [[OR_COND_I34_I_I]], label [[WHILE_COND_I33_I_I]], label [[_ZL4NANFPKC_EXIT]], !llvm.loop [[LOOP9]]
1852 // CHECK: _ZL4nanfPKc.exit:
1853 // CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I28_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_COND_I17_I_I]] ], [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I44_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_COND_I33_I_I]] ]
1854 // CHECK-NEXT: [[CONV_I:%.*]] = trunc i64 [[RETVAL_0_I_I]] to i32
1855 // CHECK-NEXT: [[BF_VALUE_I:%.*]] = and i32 [[CONV_I]], 4194303
1856 // CHECK-NEXT: [[BF_SET9_I:%.*]] = or i32 [[BF_VALUE_I]], 2143289344
1857 // CHECK-NEXT: [[TMP10:%.*]] = bitcast i32 [[BF_SET9_I]] to float
1858 // CHECK-NEXT: ret float [[TMP10]]
1860 extern "C" __device__ float test_nanf(const char *tag) {
1864 // CHECK-LABEL: @test_nan(
1865 // CHECK-NEXT: entry:
1866 // CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[TAG:%.*]], align 1, !tbaa [[TBAA3]]
1867 // CHECK-NEXT: [[CMP_I_I:%.*]] = icmp eq i8 [[TMP0]], 48
1868 // CHECK-NEXT: br i1 [[CMP_I_I]], label [[IF_THEN_I_I:%.*]], label [[WHILE_COND_I33_I_I:%.*]]
1869 // CHECK: if.then.i.i:
1870 // CHECK-NEXT: [[INCDEC_PTR_I_I:%.*]] = getelementptr inbounds i8, ptr [[TAG]], i64 1
1871 // CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr [[INCDEC_PTR_I_I]], align 1, !tbaa [[TBAA3]]
1872 // CHECK-NEXT: switch i8 [[TMP1]], label [[WHILE_COND_I17_I_I:%.*]] [
1873 // CHECK-NEXT: i8 120, label [[WHILE_COND_I_I_I_PREHEADER:%.*]]
1874 // CHECK-NEXT: i8 88, label [[WHILE_COND_I_I_I_PREHEADER]]
1876 // CHECK: while.cond.i.i.i.preheader:
1877 // CHECK-NEXT: br label [[WHILE_COND_I_I_I:%.*]]
1878 // CHECK: while.cond.i.i.i:
1879 // CHECK-NEXT: [[__TAGP_ADDR_0_I_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I_I_I:%.*]], [[CLEANUP_I_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[WHILE_COND_I_I_I_PREHEADER]] ]
1880 // CHECK-NEXT: [[__R_0_I_I_I:%.*]] = phi i64 [ [[__R_2_I_I_I:%.*]], [[CLEANUP_I_I_I]] ], [ 0, [[WHILE_COND_I_I_I_PREHEADER]] ]
1881 // CHECK-NEXT: [[TMP2:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I_I_I]], align 1, !tbaa [[TBAA3]]
1882 // CHECK-NEXT: [[CMP_NOT_I_I_I:%.*]] = icmp eq i8 [[TMP2]], 0
1883 // CHECK-NEXT: br i1 [[CMP_NOT_I_I_I]], label [[_ZL3NANPKC_EXIT:%.*]], label [[WHILE_BODY_I_I_I:%.*]]
1884 // CHECK: while.body.i.i.i:
1885 // CHECK-NEXT: [[TMP3:%.*]] = add i8 [[TMP2]], -48
1886 // CHECK-NEXT: [[OR_COND_I_I_I:%.*]] = icmp ult i8 [[TMP3]], 10
1887 // CHECK-NEXT: br i1 [[OR_COND_I_I_I]], label [[IF_END31_I_I_I:%.*]], label [[IF_ELSE_I_I_I:%.*]]
1888 // CHECK: if.else.i.i.i:
1889 // CHECK-NEXT: [[TMP4:%.*]] = add i8 [[TMP2]], -97
1890 // CHECK-NEXT: [[OR_COND33_I_I_I:%.*]] = icmp ult i8 [[TMP4]], 6
1891 // CHECK-NEXT: br i1 [[OR_COND33_I_I_I]], label [[IF_END31_I_I_I]], label [[IF_ELSE17_I_I_I:%.*]]
1892 // CHECK: if.else17.i.i.i:
1893 // CHECK-NEXT: [[TMP5:%.*]] = add i8 [[TMP2]], -65
1894 // CHECK-NEXT: [[OR_COND34_I_I_I:%.*]] = icmp ult i8 [[TMP5]], 6
1895 // CHECK-NEXT: br i1 [[OR_COND34_I_I_I]], label [[IF_END31_I_I_I]], label [[CLEANUP_I_I_I]]
1896 // CHECK: if.end31.i.i.i:
1897 // CHECK-NEXT: [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I_I_I]] ], [ -87, [[IF_ELSE_I_I_I]] ], [ -55, [[IF_ELSE17_I_I_I]] ]
1898 // CHECK-NEXT: [[MUL24_I_I_I:%.*]] = shl i64 [[__R_0_I_I_I]], 4
1899 // CHECK-NEXT: [[CONV25_I_I_I:%.*]] = sext i8 [[TMP2]] to i64
1900 // CHECK-NEXT: [[ADD26_I_I_I:%.*]] = add i64 [[MUL24_I_I_I]], [[DOTSINK]]
1901 // CHECK-NEXT: [[ADD28_I_I_I:%.*]] = add i64 [[ADD26_I_I_I]], [[CONV25_I_I_I]]
1902 // CHECK-NEXT: [[INCDEC_PTR_I_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I_I_I]], i64 1
1903 // CHECK-NEXT: br label [[CLEANUP_I_I_I]]
1904 // CHECK: cleanup.i.i.i:
1905 // CHECK-NEXT: [[__TAGP_ADDR_1_I_I_I]] = phi ptr [ [[INCDEC_PTR_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__TAGP_ADDR_0_I_I_I]], [[IF_ELSE17_I_I_I]] ]
1906 // CHECK-NEXT: [[__R_2_I_I_I]] = phi i64 [ [[ADD28_I_I_I]], [[IF_END31_I_I_I]] ], [ [[__R_0_I_I_I]], [[IF_ELSE17_I_I_I]] ]
1907 // CHECK-NEXT: [[COND_I_I_I:%.*]] = phi i1 [ true, [[IF_END31_I_I_I]] ], [ false, [[IF_ELSE17_I_I_I]] ]
1908 // CHECK-NEXT: br i1 [[COND_I_I_I]], label [[WHILE_COND_I_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP10]]
1909 // CHECK: while.cond.i17.i.i:
1910 // CHECK-NEXT: [[__TAGP_ADDR_0_I14_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I26_I_I:%.*]], [[CLEANUP_I28_I_I:%.*]] ], [ [[INCDEC_PTR_I_I]], [[IF_THEN_I_I]] ]
1911 // CHECK-NEXT: [[__R_0_I15_I_I:%.*]] = phi i64 [ [[__R_1_I27_I_I:%.*]], [[CLEANUP_I28_I_I]] ], [ 0, [[IF_THEN_I_I]] ]
1912 // CHECK-NEXT: [[TMP6:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I14_I_I]], align 1, !tbaa [[TBAA3]]
1913 // CHECK-NEXT: [[CMP_NOT_I16_I_I:%.*]] = icmp eq i8 [[TMP6]], 0
1914 // CHECK-NEXT: br i1 [[CMP_NOT_I16_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I19_I_I:%.*]]
1915 // CHECK: while.body.i19.i.i:
1916 // CHECK-NEXT: [[TMP7:%.*]] = and i8 [[TMP6]], -8
1917 // CHECK-NEXT: [[OR_COND_I18_I_I:%.*]] = icmp eq i8 [[TMP7]], 48
1918 // CHECK-NEXT: br i1 [[OR_COND_I18_I_I]], label [[IF_THEN_I25_I_I:%.*]], label [[CLEANUP_I28_I_I]]
1919 // CHECK: if.then.i25.i.i:
1920 // CHECK-NEXT: [[MUL_I20_I_I:%.*]] = shl i64 [[__R_0_I15_I_I]], 3
1921 // CHECK-NEXT: [[CONV5_I21_I_I:%.*]] = sext i8 [[TMP6]] to i64
1922 // CHECK-NEXT: [[ADD_I22_I_I:%.*]] = add i64 [[MUL_I20_I_I]], -48
1923 // CHECK-NEXT: [[SUB_I23_I_I:%.*]] = add i64 [[ADD_I22_I_I]], [[CONV5_I21_I_I]]
1924 // CHECK-NEXT: [[INCDEC_PTR_I24_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I14_I_I]], i64 1
1925 // CHECK-NEXT: br label [[CLEANUP_I28_I_I]]
1926 // CHECK: cleanup.i28.i.i:
1927 // CHECK-NEXT: [[__TAGP_ADDR_1_I26_I_I]] = phi ptr [ [[INCDEC_PTR_I24_I_I]], [[IF_THEN_I25_I_I]] ], [ [[__TAGP_ADDR_0_I14_I_I]], [[WHILE_BODY_I19_I_I]] ]
1928 // CHECK-NEXT: [[__R_1_I27_I_I]] = phi i64 [ [[SUB_I23_I_I]], [[IF_THEN_I25_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_BODY_I19_I_I]] ]
1929 // CHECK-NEXT: br i1 [[OR_COND_I18_I_I]], label [[WHILE_COND_I17_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP6]]
1930 // CHECK: while.cond.i33.i.i:
1931 // CHECK-NEXT: [[__TAGP_ADDR_0_I30_I_I:%.*]] = phi ptr [ [[__TAGP_ADDR_1_I42_I_I:%.*]], [[CLEANUP_I44_I_I:%.*]] ], [ [[TAG]], [[ENTRY:%.*]] ]
1932 // CHECK-NEXT: [[__R_0_I31_I_I:%.*]] = phi i64 [ [[__R_1_I43_I_I:%.*]], [[CLEANUP_I44_I_I]] ], [ 0, [[ENTRY]] ]
1933 // CHECK-NEXT: [[TMP8:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I30_I_I]], align 1, !tbaa [[TBAA3]]
1934 // CHECK-NEXT: [[CMP_NOT_I32_I_I:%.*]] = icmp eq i8 [[TMP8]], 0
1935 // CHECK-NEXT: br i1 [[CMP_NOT_I32_I_I]], label [[_ZL3NANPKC_EXIT]], label [[WHILE_BODY_I35_I_I:%.*]]
1936 // CHECK: while.body.i35.i.i:
1937 // CHECK-NEXT: [[TMP9:%.*]] = add i8 [[TMP8]], -48
1938 // CHECK-NEXT: [[OR_COND_I34_I_I:%.*]] = icmp ult i8 [[TMP9]], 10
1939 // CHECK-NEXT: br i1 [[OR_COND_I34_I_I]], label [[IF_THEN_I41_I_I:%.*]], label [[CLEANUP_I44_I_I]]
1940 // CHECK: if.then.i41.i.i:
1941 // CHECK-NEXT: [[MUL_I36_I_I:%.*]] = mul i64 [[__R_0_I31_I_I]], 10
1942 // CHECK-NEXT: [[CONV5_I37_I_I:%.*]] = sext i8 [[TMP8]] to i64
1943 // CHECK-NEXT: [[ADD_I38_I_I:%.*]] = add i64 [[MUL_I36_I_I]], -48
1944 // CHECK-NEXT: [[SUB_I39_I_I:%.*]] = add i64 [[ADD_I38_I_I]], [[CONV5_I37_I_I]]
1945 // CHECK-NEXT: [[INCDEC_PTR_I40_I_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I30_I_I]], i64 1
1946 // CHECK-NEXT: br label [[CLEANUP_I44_I_I]]
1947 // CHECK: cleanup.i44.i.i:
1948 // CHECK-NEXT: [[__TAGP_ADDR_1_I42_I_I]] = phi ptr [ [[INCDEC_PTR_I40_I_I]], [[IF_THEN_I41_I_I]] ], [ [[__TAGP_ADDR_0_I30_I_I]], [[WHILE_BODY_I35_I_I]] ]
1949 // CHECK-NEXT: [[__R_1_I43_I_I]] = phi i64 [ [[SUB_I39_I_I]], [[IF_THEN_I41_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_BODY_I35_I_I]] ]
1950 // CHECK-NEXT: br i1 [[OR_COND_I34_I_I]], label [[WHILE_COND_I33_I_I]], label [[_ZL3NANPKC_EXIT]], !llvm.loop [[LOOP9]]
1951 // CHECK: _ZL3nanPKc.exit:
1952 // CHECK-NEXT: [[RETVAL_0_I_I:%.*]] = phi i64 [ 0, [[CLEANUP_I28_I_I]] ], [ [[__R_0_I15_I_I]], [[WHILE_COND_I17_I_I]] ], [ 0, [[CLEANUP_I_I_I]] ], [ [[__R_0_I_I_I]], [[WHILE_COND_I_I_I]] ], [ 0, [[CLEANUP_I44_I_I]] ], [ [[__R_0_I31_I_I]], [[WHILE_COND_I33_I_I]] ]
1953 // CHECK-NEXT: [[BF_VALUE_I:%.*]] = and i64 [[RETVAL_0_I_I]], 2251799813685247
1954 // CHECK-NEXT: [[BF_SET9_I:%.*]] = or i64 [[BF_VALUE_I]], 9221120237041090560
1955 // CHECK-NEXT: [[TMP10:%.*]] = bitcast i64 [[BF_SET9_I]] to double
1956 // CHECK-NEXT: ret double [[TMP10]]
1958 extern "C" __device__ double test_nan(const char *tag) {
1962 // CHECK-LABEL: @test_nanf_emptystr(
1963 // CHECK-NEXT: entry:
1964 // CHECK-NEXT: ret float 0x7FF8000000000000
1966 extern "C" __device__ float test_nanf_emptystr() {
1970 // CHECK-LABEL: @test_nan_emptystr(
1971 // CHECK-NEXT: entry:
1972 // CHECK-NEXT: ret double 0x7FF8000000000000
1974 extern "C" __device__ double test_nan_emptystr() {
1978 // CHECK-LABEL: @test_nanf_fill(
1979 // CHECK-NEXT: entry:
1980 // CHECK-NEXT: ret float 0x7FF8000000000000
1982 extern "C" __device__ float test_nanf_fill() {
1983 return nanf("0x456");
1986 // CHECK-LABEL: @test_nan_fill(
1987 // CHECK-NEXT: entry:
1988 // CHECK-NEXT: ret double 0x7FF8000000000000
1990 extern "C" __device__ double test_nan_fill() {
1991 return nan("0x123");
1994 // DEFAULT-LABEL: @test_nearbyintf(
1995 // DEFAULT-NEXT: entry:
1996 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_nearbyint_f32(float noundef [[X:%.*]]) #[[ATTR13]]
1997 // DEFAULT-NEXT: ret float [[CALL_I]]
1999 // FINITEONLY-LABEL: @test_nearbyintf(
2000 // FINITEONLY-NEXT: entry:
2001 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_nearbyint_f32(float noundef [[X:%.*]]) #[[ATTR13]]
2002 // FINITEONLY-NEXT: ret float [[CALL_I]]
2004 extern "C" __device__ float test_nearbyintf(float x) {
2005 return nearbyintf(x);
2008 // DEFAULT-LABEL: @test_nearbyint(
2009 // DEFAULT-NEXT: entry:
2010 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_nearbyint_f64(double noundef [[X:%.*]]) #[[ATTR13]]
2011 // DEFAULT-NEXT: ret double [[CALL_I]]
2013 // FINITEONLY-LABEL: @test_nearbyint(
2014 // FINITEONLY-NEXT: entry:
2015 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_nearbyint_f64(double noundef [[X:%.*]]) #[[ATTR13]]
2016 // FINITEONLY-NEXT: ret double [[CALL_I]]
2018 extern "C" __device__ double test_nearbyint(double x) {
2019 return nearbyint(x);
2022 // DEFAULT-LABEL: @test_nextafterf(
2023 // DEFAULT-NEXT: entry:
2024 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_nextafter_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
2025 // DEFAULT-NEXT: ret float [[CALL_I]]
2027 // FINITEONLY-LABEL: @test_nextafterf(
2028 // FINITEONLY-NEXT: entry:
2029 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_nextafter_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
2030 // FINITEONLY-NEXT: ret float [[CALL_I]]
2032 extern "C" __device__ float test_nextafterf(float x, float y) {
2033 return nextafterf(x, y);
2036 // DEFAULT-LABEL: @test_nextafter(
2037 // DEFAULT-NEXT: entry:
2038 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_nextafter_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
2039 // DEFAULT-NEXT: ret double [[CALL_I]]
2041 // FINITEONLY-LABEL: @test_nextafter(
2042 // FINITEONLY-NEXT: entry:
2043 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_nextafter_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
2044 // FINITEONLY-NEXT: ret double [[CALL_I]]
2046 extern "C" __device__ double test_nextafter(double x, double y) {
2047 return nextafter(x, y);
2050 // DEFAULT-LABEL: @test_norm3df(
2051 // DEFAULT-NEXT: entry:
2052 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_len3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR13]]
2053 // DEFAULT-NEXT: ret float [[CALL_I]]
2055 // FINITEONLY-LABEL: @test_norm3df(
2056 // FINITEONLY-NEXT: entry:
2057 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_len3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR13]]
2058 // FINITEONLY-NEXT: ret float [[CALL_I]]
2060 extern "C" __device__ float test_norm3df(float x, float y, float z) {
2061 return norm3df(x, y, z);
2064 // DEFAULT-LABEL: @test_norm3d(
2065 // DEFAULT-NEXT: entry:
2066 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_len3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR13]]
2067 // DEFAULT-NEXT: ret double [[CALL_I]]
2069 // FINITEONLY-LABEL: @test_norm3d(
2070 // FINITEONLY-NEXT: entry:
2071 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_len3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR13]]
2072 // FINITEONLY-NEXT: ret double [[CALL_I]]
2074 extern "C" __device__ double test_norm3d(double x, double y, double z) {
2075 return norm3d(x, y, z);
2078 // DEFAULT-LABEL: @test_norm4df(
2079 // DEFAULT-NEXT: entry:
2080 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_len4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR13]]
2081 // DEFAULT-NEXT: ret float [[CALL_I]]
2083 // FINITEONLY-LABEL: @test_norm4df(
2084 // FINITEONLY-NEXT: entry:
2085 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_len4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR13]]
2086 // FINITEONLY-NEXT: ret float [[CALL_I]]
2088 extern "C" __device__ float test_norm4df(float x, float y, float z, float w) {
2089 return norm4df(x, y, z, w);
2092 // DEFAULT-LABEL: @test_norm4d(
2093 // DEFAULT-NEXT: entry:
2094 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_len4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR13]]
2095 // DEFAULT-NEXT: ret double [[CALL_I]]
2097 // FINITEONLY-LABEL: @test_norm4d(
2098 // FINITEONLY-NEXT: entry:
2099 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_len4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR13]]
2100 // FINITEONLY-NEXT: ret double [[CALL_I]]
2102 extern "C" __device__ double test_norm4d(double x, double y, double z, double w) {
2103 return norm4d(x, y, z, w);
2106 // DEFAULT-LABEL: @test_normcdff(
2107 // DEFAULT-NEXT: entry:
2108 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ncdf_f32(float noundef [[X:%.*]]) #[[ATTR14]]
2109 // DEFAULT-NEXT: ret float [[CALL_I]]
2111 // FINITEONLY-LABEL: @test_normcdff(
2112 // FINITEONLY-NEXT: entry:
2113 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ncdf_f32(float noundef [[X:%.*]]) #[[ATTR14]]
2114 // FINITEONLY-NEXT: ret float [[CALL_I]]
2116 extern "C" __device__ float test_normcdff(float x) {
2120 // DEFAULT-LABEL: @test_normcdf(
2121 // DEFAULT-NEXT: entry:
2122 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ncdf_f64(double noundef [[X:%.*]]) #[[ATTR14]]
2123 // DEFAULT-NEXT: ret double [[CALL_I]]
2125 // FINITEONLY-LABEL: @test_normcdf(
2126 // FINITEONLY-NEXT: entry:
2127 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ncdf_f64(double noundef [[X:%.*]]) #[[ATTR14]]
2128 // FINITEONLY-NEXT: ret double [[CALL_I]]
2130 extern "C" __device__ double test_normcdf(double x) {
2134 // DEFAULT-LABEL: @test_normcdfinvf(
2135 // DEFAULT-NEXT: entry:
2136 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_ncdfinv_f32(float noundef [[X:%.*]]) #[[ATTR14]]
2137 // DEFAULT-NEXT: ret float [[CALL_I]]
2139 // FINITEONLY-LABEL: @test_normcdfinvf(
2140 // FINITEONLY-NEXT: entry:
2141 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ncdfinv_f32(float noundef [[X:%.*]]) #[[ATTR14]]
2142 // FINITEONLY-NEXT: ret float [[CALL_I]]
2144 extern "C" __device__ float test_normcdfinvf(float x) {
2145 return normcdfinvf(x);
2148 // DEFAULT-LABEL: @test_normcdfinv(
2149 // DEFAULT-NEXT: entry:
2150 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_ncdfinv_f64(double noundef [[X:%.*]]) #[[ATTR14]]
2151 // DEFAULT-NEXT: ret double [[CALL_I]]
2153 // FINITEONLY-LABEL: @test_normcdfinv(
2154 // FINITEONLY-NEXT: entry:
2155 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ncdfinv_f64(double noundef [[X:%.*]]) #[[ATTR14]]
2156 // FINITEONLY-NEXT: ret double [[CALL_I]]
2158 extern "C" __device__ double test_normcdfinv(double x) {
2159 return normcdfinv(x);
2162 // DEFAULT-LABEL: @test_normf(
2163 // DEFAULT-NEXT: entry:
2164 // DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
2165 // DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
2166 // DEFAULT: while.body.i:
2167 // DEFAULT-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
2168 // DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
2169 // DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
2170 // DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
2171 // DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA15]]
2172 // DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]]
2173 // DEFAULT-NEXT: [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]]
2174 // DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
2175 // DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
2176 // DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP19:![0-9]+]]
2177 // DEFAULT: _ZL5normfiPKf.exit:
2178 // DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
2179 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR13]]
2180 // DEFAULT-NEXT: ret float [[CALL_I]]
2182 // FINITEONLY-LABEL: @test_normf(
2183 // FINITEONLY-NEXT: entry:
2184 // FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
2185 // FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
2186 // FINITEONLY: while.body.i:
2187 // FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
2188 // FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
2189 // FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
2190 // FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
2191 // FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA15]]
2192 // FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[TMP0]], [[TMP0]]
2193 // FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract float [[__R_08_I]], [[MUL_I]]
2194 // FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
2195 // FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
2196 // FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP19:![0-9]+]]
2197 // FINITEONLY: _ZL5normfiPKf.exit:
2198 // FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
2199 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR13]]
2200 // FINITEONLY-NEXT: ret float [[CALL_I]]
2202 extern "C" __device__ float test_normf(int x, const float *y) {
2206 // DEFAULT-LABEL: @test_norm(
2207 // DEFAULT-NEXT: entry:
2208 // DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
2209 // DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
2210 // DEFAULT: while.body.i:
2211 // DEFAULT-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
2212 // DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
2213 // DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
2214 // DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
2215 // DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA17]]
2216 // DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]]
2217 // DEFAULT-NEXT: [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]]
2218 // DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
2219 // DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
2220 // DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
2221 // DEFAULT: _ZL4normiPKd.exit:
2222 // DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
2223 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR13]]
2224 // DEFAULT-NEXT: ret double [[CALL_I]]
2226 // FINITEONLY-LABEL: @test_norm(
2227 // FINITEONLY-NEXT: entry:
2228 // FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
2229 // FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
2230 // FINITEONLY: while.body.i:
2231 // FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
2232 // FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
2233 // FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
2234 // FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
2235 // FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA17]]
2236 // FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract double [[TMP0]], [[TMP0]]
2237 // FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract double [[__R_08_I]], [[MUL_I]]
2238 // FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
2239 // FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
2240 // FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
2241 // FINITEONLY: _ZL4normiPKd.exit:
2242 // FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
2243 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR13]]
2244 // FINITEONLY-NEXT: ret double [[CALL_I]]
2246 extern "C" __device__ double test_norm(int x, const double *y) {
2250 // DEFAULT-LABEL: @test_powf(
2251 // DEFAULT-NEXT: entry:
2252 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
2253 // DEFAULT-NEXT: ret float [[CALL_I]]
2255 // FINITEONLY-LABEL: @test_powf(
2256 // FINITEONLY-NEXT: entry:
2257 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
2258 // FINITEONLY-NEXT: ret float [[CALL_I]]
2260 extern "C" __device__ float test_powf(float x, float y) {
2264 // DEFAULT-LABEL: @test_pow(
2265 // DEFAULT-NEXT: entry:
2266 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_pow_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
2267 // DEFAULT-NEXT: ret double [[CALL_I]]
2269 // FINITEONLY-LABEL: @test_pow(
2270 // FINITEONLY-NEXT: entry:
2271 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_pow_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
2272 // FINITEONLY-NEXT: ret double [[CALL_I]]
2274 extern "C" __device__ double test_pow(double x, double y) {
2278 // DEFAULT-LABEL: @test_powif(
2279 // DEFAULT-NEXT: entry:
2280 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_pown_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
2281 // DEFAULT-NEXT: ret float [[CALL_I]]
2283 // FINITEONLY-LABEL: @test_powif(
2284 // FINITEONLY-NEXT: entry:
2285 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_pown_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
2286 // FINITEONLY-NEXT: ret float [[CALL_I]]
2288 extern "C" __device__ float test_powif(float x, int y) {
2292 // DEFAULT-LABEL: @test_powi(
2293 // DEFAULT-NEXT: entry:
2294 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_pown_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
2295 // DEFAULT-NEXT: ret double [[CALL_I]]
2297 // FINITEONLY-LABEL: @test_powi(
2298 // FINITEONLY-NEXT: entry:
2299 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_pown_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
2300 // FINITEONLY-NEXT: ret double [[CALL_I]]
2302 extern "C" __device__ double test_powi(double x, int y) {
2306 // DEFAULT-LABEL: @test_rcbrtf(
2307 // DEFAULT-NEXT: entry:
2308 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rcbrt_f32(float noundef [[X:%.*]]) #[[ATTR14]]
2309 // DEFAULT-NEXT: ret float [[CALL_I]]
2311 // FINITEONLY-LABEL: @test_rcbrtf(
2312 // FINITEONLY-NEXT: entry:
2313 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rcbrt_f32(float noundef [[X:%.*]]) #[[ATTR14]]
2314 // FINITEONLY-NEXT: ret float [[CALL_I]]
2316 extern "C" __device__ float test_rcbrtf(float x) {
2320 // DEFAULT-LABEL: @test_rcbrt(
2321 // DEFAULT-NEXT: entry:
2322 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rcbrt_f64(double noundef [[X:%.*]]) #[[ATTR14]]
2323 // DEFAULT-NEXT: ret double [[CALL_I]]
2325 // FINITEONLY-LABEL: @test_rcbrt(
2326 // FINITEONLY-NEXT: entry:
2327 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rcbrt_f64(double noundef [[X:%.*]]) #[[ATTR14]]
2328 // FINITEONLY-NEXT: ret double [[CALL_I]]
2330 extern "C" __device__ double test_rcbrt(double x) {
2334 // DEFAULT-LABEL: @test_remainderf(
2335 // DEFAULT-NEXT: entry:
2336 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_remainder_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
2337 // DEFAULT-NEXT: ret float [[CALL_I]]
2339 // FINITEONLY-LABEL: @test_remainderf(
2340 // FINITEONLY-NEXT: entry:
2341 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_remainder_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
2342 // FINITEONLY-NEXT: ret float [[CALL_I]]
2344 extern "C" __device__ float test_remainderf(float x, float y) {
2345 return remainderf(x, y);
2348 // DEFAULT-LABEL: @test_remainder(
2349 // DEFAULT-NEXT: entry:
2350 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_remainder_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
2351 // DEFAULT-NEXT: ret double [[CALL_I]]
2353 // FINITEONLY-LABEL: @test_remainder(
2354 // FINITEONLY-NEXT: entry:
2355 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_remainder_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
2356 // FINITEONLY-NEXT: ret double [[CALL_I]]
2358 extern "C" __device__ double test_remainder(double x, double y) {
2359 return remainder(x, y);
2362 // DEFAULT-LABEL: @test_remquof(
2363 // DEFAULT-NEXT: entry:
2364 // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
2365 // DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2366 // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
2367 // DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]]
2368 // DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA11]]
2369 // DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2370 // DEFAULT-NEXT: ret float [[CALL_I]]
2372 // FINITEONLY-LABEL: @test_remquof(
2373 // FINITEONLY-NEXT: entry:
2374 // FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
2375 // FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2376 // FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
2377 // FINITEONLY-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]]
2378 // FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA11]]
2379 // FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2380 // FINITEONLY-NEXT: ret float [[CALL_I]]
2382 extern "C" __device__ float test_remquof(float x, float y, int* z) {
2383 return remquof(x, y, z);
2386 // DEFAULT-LABEL: @test_remquo(
2387 // DEFAULT-NEXT: entry:
2388 // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
2389 // DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2390 // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
2391 // DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]]
2392 // DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA11]]
2393 // DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2394 // DEFAULT-NEXT: ret double [[CALL_I]]
2396 // FINITEONLY-LABEL: @test_remquo(
2397 // FINITEONLY-NEXT: entry:
2398 // FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
2399 // FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2400 // FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
2401 // FINITEONLY-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA11]]
2402 // FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA11]]
2403 // FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2404 // FINITEONLY-NEXT: ret double [[CALL_I]]
2406 extern "C" __device__ double test_remquo(double x, double y, int* z) {
2407 return remquo(x, y, z);
2410 // DEFAULT-LABEL: @test_rhypotf(
2411 // DEFAULT-NEXT: entry:
2412 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rhypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
2413 // DEFAULT-NEXT: ret float [[CALL_I]]
2415 // FINITEONLY-LABEL: @test_rhypotf(
2416 // FINITEONLY-NEXT: entry:
2417 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rhypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
2418 // FINITEONLY-NEXT: ret float [[CALL_I]]
2420 extern "C" __device__ float test_rhypotf(float x, float y) {
2421 return rhypotf(x, y);
2424 // DEFAULT-LABEL: @test_rhypot(
2425 // DEFAULT-NEXT: entry:
2426 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rhypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
2427 // DEFAULT-NEXT: ret double [[CALL_I]]
2429 // FINITEONLY-LABEL: @test_rhypot(
2430 // FINITEONLY-NEXT: entry:
2431 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rhypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
2432 // FINITEONLY-NEXT: ret double [[CALL_I]]
2434 extern "C" __device__ double test_rhypot(double x, double y) {
2435 return rhypot(x, y);
2438 // DEFAULT-LABEL: @test_rintf(
2439 // DEFAULT-NEXT: entry:
2440 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR13]]
2441 // DEFAULT-NEXT: ret float [[CALL_I]]
2443 // FINITEONLY-LABEL: @test_rintf(
2444 // FINITEONLY-NEXT: entry:
2445 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR13]]
2446 // FINITEONLY-NEXT: ret float [[CALL_I]]
2448 extern "C" __device__ float test_rintf(float x) {
2452 // DEFAULT-LABEL: @test_rint(
2453 // DEFAULT-NEXT: entry:
2454 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR13]]
2455 // DEFAULT-NEXT: ret double [[CALL_I]]
2457 // FINITEONLY-LABEL: @test_rint(
2458 // FINITEONLY-NEXT: entry:
2459 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR13]]
2460 // FINITEONLY-NEXT: ret double [[CALL_I]]
2462 extern "C" __device__ double test_rint(double x) {
2466 // DEFAULT-LABEL: @test_rnormf(
2467 // DEFAULT-NEXT: entry:
2468 // DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
2469 // DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
2470 // DEFAULT: while.body.i:
2471 // DEFAULT-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
2472 // DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
2473 // DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
2474 // DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
2475 // DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA15]]
2476 // DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]]
2477 // DEFAULT-NEXT: [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]]
2478 // DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
2479 // DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
2480 // DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
2481 // DEFAULT: _ZL6rnormfiPKf.exit:
2482 // DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
2483 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
2484 // DEFAULT-NEXT: ret float [[CALL_I]]
2486 // FINITEONLY-LABEL: @test_rnormf(
2487 // FINITEONLY-NEXT: entry:
2488 // FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
2489 // FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
2490 // FINITEONLY: while.body.i:
2491 // FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
2492 // FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
2493 // FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
2494 // FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
2495 // FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA15]]
2496 // FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[TMP0]], [[TMP0]]
2497 // FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract float [[__R_08_I]], [[MUL_I]]
2498 // FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
2499 // FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
2500 // FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
2501 // FINITEONLY: _ZL6rnormfiPKf.exit:
2502 // FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
2503 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
2504 // FINITEONLY-NEXT: ret float [[CALL_I]]
2506 extern "C" __device__ float test_rnormf(int x, const float* y) {
2507 return rnormf(x, y);
2510 // DEFAULT-LABEL: @test_rnorm(
2511 // DEFAULT-NEXT: entry:
2512 // DEFAULT-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
2513 // DEFAULT-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
2514 // DEFAULT: while.body.i:
2515 // DEFAULT-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
2516 // DEFAULT-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
2517 // DEFAULT-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
2518 // DEFAULT-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
2519 // DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA17]]
2520 // DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]]
2521 // DEFAULT-NEXT: [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]]
2522 // DEFAULT-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
2523 // DEFAULT-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
2524 // DEFAULT-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]]
2525 // DEFAULT: _ZL5rnormiPKd.exit:
2526 // DEFAULT-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
2527 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
2528 // DEFAULT-NEXT: ret double [[CALL_I]]
2530 // FINITEONLY-LABEL: @test_rnorm(
2531 // FINITEONLY-NEXT: entry:
2532 // FINITEONLY-NEXT: [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
2533 // FINITEONLY-NEXT: br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
2534 // FINITEONLY: while.body.i:
2535 // FINITEONLY-NEXT: [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
2536 // FINITEONLY-NEXT: [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
2537 // FINITEONLY-NEXT: [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
2538 // FINITEONLY-NEXT: [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
2539 // FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA17]]
2540 // FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract double [[TMP0]], [[TMP0]]
2541 // FINITEONLY-NEXT: [[ADD_I]] = fadd nnan ninf contract double [[__R_08_I]], [[MUL_I]]
2542 // FINITEONLY-NEXT: [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
2543 // FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
2544 // FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]]
2545 // FINITEONLY: _ZL5rnormiPKd.exit:
2546 // FINITEONLY-NEXT: [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
2547 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
2548 // FINITEONLY-NEXT: ret double [[CALL_I]]
2550 extern "C" __device__ double test_rnorm(int x, const double* y) {
2554 // DEFAULT-LABEL: @test_rnorm3df(
2555 // DEFAULT-NEXT: entry:
2556 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rlen3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR13]]
2557 // DEFAULT-NEXT: ret float [[CALL_I]]
2559 // FINITEONLY-LABEL: @test_rnorm3df(
2560 // FINITEONLY-NEXT: entry:
2561 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rlen3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR13]]
2562 // FINITEONLY-NEXT: ret float [[CALL_I]]
2564 extern "C" __device__ float test_rnorm3df(float x, float y, float z) {
2565 return rnorm3df(x, y, z);
2568 // DEFAULT-LABEL: @test_rnorm3d(
2569 // DEFAULT-NEXT: entry:
2570 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rlen3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR13]]
2571 // DEFAULT-NEXT: ret double [[CALL_I]]
2573 // FINITEONLY-LABEL: @test_rnorm3d(
2574 // FINITEONLY-NEXT: entry:
2575 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rlen3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR13]]
2576 // FINITEONLY-NEXT: ret double [[CALL_I]]
2578 extern "C" __device__ double test_rnorm3d(double x, double y, double z) {
2579 return rnorm3d(x, y, z);
2582 // DEFAULT-LABEL: @test_rnorm4df(
2583 // DEFAULT-NEXT: entry:
2584 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rlen4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR13]]
2585 // DEFAULT-NEXT: ret float [[CALL_I]]
2587 // FINITEONLY-LABEL: @test_rnorm4df(
2588 // FINITEONLY-NEXT: entry:
2589 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rlen4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR13]]
2590 // FINITEONLY-NEXT: ret float [[CALL_I]]
2592 extern "C" __device__ float test_rnorm4df(float x, float y, float z, float w) {
2593 return rnorm4df(x, y, z, w);
2596 // DEFAULT-LABEL: @test_rnorm4d(
2597 // DEFAULT-NEXT: entry:
2598 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rlen4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR13]]
2599 // DEFAULT-NEXT: ret double [[CALL_I]]
2601 // FINITEONLY-LABEL: @test_rnorm4d(
2602 // FINITEONLY-NEXT: entry:
2603 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rlen4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR13]]
2604 // FINITEONLY-NEXT: ret double [[CALL_I]]
2606 extern "C" __device__ double test_rnorm4d(double x, double y, double z, double w) {
2607 return rnorm4d(x, y, z, w);
2610 // DEFAULT-LABEL: @test_roundf(
2611 // DEFAULT-NEXT: entry:
2612 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR13]]
2613 // DEFAULT-NEXT: ret float [[CALL_I]]
2615 // FINITEONLY-LABEL: @test_roundf(
2616 // FINITEONLY-NEXT: entry:
2617 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR13]]
2618 // FINITEONLY-NEXT: ret float [[CALL_I]]
2620 extern "C" __device__ float test_roundf(float x) {
2624 // DEFAULT-LABEL: @test_round(
2625 // DEFAULT-NEXT: entry:
2626 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR13]]
2627 // DEFAULT-NEXT: ret double [[CALL_I]]
2629 // FINITEONLY-LABEL: @test_round(
2630 // FINITEONLY-NEXT: entry:
2631 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR13]]
2632 // FINITEONLY-NEXT: ret double [[CALL_I]]
2634 extern "C" __device__ double test_round(double x) {
2638 // DEFAULT-LABEL: @test_rsqrtf(
2639 // DEFAULT-NEXT: entry:
2640 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]]
2641 // DEFAULT-NEXT: ret float [[CALL_I]]
2643 // FINITEONLY-LABEL: @test_rsqrtf(
2644 // FINITEONLY-NEXT: entry:
2645 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rsqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]]
2646 // FINITEONLY-NEXT: ret float [[CALL_I]]
2648 extern "C" __device__ float test_rsqrtf(float x) {
2652 // DEFAULT-LABEL: @test_rsqrt(
2653 // DEFAULT-NEXT: entry:
2654 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]]
2655 // DEFAULT-NEXT: ret double [[CALL_I]]
2657 // FINITEONLY-LABEL: @test_rsqrt(
2658 // FINITEONLY-NEXT: entry:
2659 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rsqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]]
2660 // FINITEONLY-NEXT: ret double [[CALL_I]]
2662 extern "C" __device__ double test_rsqrt(double x) {
2666 // DEFAULT-LABEL: @test_scalblnf(
2667 // DEFAULT-NEXT: entry:
2668 // DEFAULT-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i64 [[Y:%.*]], 9223372036854775807
2669 // DEFAULT-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]]
2670 // DEFAULT: cond.true.i:
2671 // DEFAULT-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32
2672 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR13]]
2673 // DEFAULT-NEXT: br label [[_ZL8SCALBLNFFL_EXIT:%.*]]
2674 // DEFAULT: cond.false.i:
2675 // DEFAULT-NEXT: [[CALL2_I:%.*]] = tail call contract float @__ocml_scalb_f32(float noundef [[X]], float noundef 0x43E0000000000000) #[[ATTR13]]
2676 // DEFAULT-NEXT: br label [[_ZL8SCALBLNFFL_EXIT]]
2677 // DEFAULT: _ZL8scalblnffl.exit:
2678 // DEFAULT-NEXT: [[COND_I:%.*]] = phi contract float [ [[CALL_I]], [[COND_TRUE_I]] ], [ [[CALL2_I]], [[COND_FALSE_I]] ]
2679 // DEFAULT-NEXT: ret float [[COND_I]]
2681 // FINITEONLY-LABEL: @test_scalblnf(
2682 // FINITEONLY-NEXT: entry:
2683 // FINITEONLY-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i64 [[Y:%.*]], 9223372036854775807
2684 // FINITEONLY-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]]
2685 // FINITEONLY: cond.true.i:
2686 // FINITEONLY-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32
2687 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR13]]
2688 // FINITEONLY-NEXT: br label [[_ZL8SCALBLNFFL_EXIT:%.*]]
2689 // FINITEONLY: cond.false.i:
2690 // FINITEONLY-NEXT: [[CALL2_I:%.*]] = tail call nnan ninf contract float @__ocml_scalb_f32(float noundef [[X]], float noundef 0x43E0000000000000) #[[ATTR13]]
2691 // FINITEONLY-NEXT: br label [[_ZL8SCALBLNFFL_EXIT]]
2692 // FINITEONLY: _ZL8scalblnffl.exit:
2693 // FINITEONLY-NEXT: [[COND_I:%.*]] = phi nnan ninf contract float [ [[CALL_I]], [[COND_TRUE_I]] ], [ [[CALL2_I]], [[COND_FALSE_I]] ]
2694 // FINITEONLY-NEXT: ret float [[COND_I]]
2696 extern "C" __device__ float test_scalblnf(float x, long int y) {
2697 return scalblnf(x, y);
2700 // DEFAULT-LABEL: @test_scalbln(
2701 // DEFAULT-NEXT: entry:
2702 // DEFAULT-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i64 [[Y:%.*]], 9223372036854775807
2703 // DEFAULT-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]]
2704 // DEFAULT: cond.true.i:
2705 // DEFAULT-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32
2706 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR13]]
2707 // DEFAULT-NEXT: br label [[_ZL7SCALBLNDL_EXIT:%.*]]
2708 // DEFAULT: cond.false.i:
2709 // DEFAULT-NEXT: [[CALL2_I:%.*]] = tail call contract double @__ocml_scalb_f64(double noundef [[X]], double noundef 0x43E0000000000000) #[[ATTR13]]
2710 // DEFAULT-NEXT: br label [[_ZL7SCALBLNDL_EXIT]]
2711 // DEFAULT: _ZL7scalblndl.exit:
2712 // DEFAULT-NEXT: [[COND_I:%.*]] = phi contract double [ [[CALL_I]], [[COND_TRUE_I]] ], [ [[CALL2_I]], [[COND_FALSE_I]] ]
2713 // DEFAULT-NEXT: ret double [[COND_I]]
2715 // FINITEONLY-LABEL: @test_scalbln(
2716 // FINITEONLY-NEXT: entry:
2717 // FINITEONLY-NEXT: [[CMP_NOT_I:%.*]] = icmp eq i64 [[Y:%.*]], 9223372036854775807
2718 // FINITEONLY-NEXT: br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]]
2719 // FINITEONLY: cond.true.i:
2720 // FINITEONLY-NEXT: [[CONV_I:%.*]] = trunc i64 [[Y]] to i32
2721 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR13]]
2722 // FINITEONLY-NEXT: br label [[_ZL7SCALBLNDL_EXIT:%.*]]
2723 // FINITEONLY: cond.false.i:
2724 // FINITEONLY-NEXT: [[CALL2_I:%.*]] = tail call nnan ninf contract double @__ocml_scalb_f64(double noundef [[X]], double noundef 0x43E0000000000000) #[[ATTR13]]
2725 // FINITEONLY-NEXT: br label [[_ZL7SCALBLNDL_EXIT]]
2726 // FINITEONLY: _ZL7scalblndl.exit:
2727 // FINITEONLY-NEXT: [[COND_I:%.*]] = phi nnan ninf contract double [ [[CALL_I]], [[COND_TRUE_I]] ], [ [[CALL2_I]], [[COND_FALSE_I]] ]
2728 // FINITEONLY-NEXT: ret double [[COND_I]]
2730 extern "C" __device__ double test_scalbln(double x, long int y) {
2731 return scalbln(x, y);
2734 // DEFAULT-LABEL: @test_scalbnf(
2735 // DEFAULT-NEXT: entry:
2736 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]]
2737 // DEFAULT-NEXT: ret float [[CALL_I]]
2739 // FINITEONLY-LABEL: @test_scalbnf(
2740 // FINITEONLY-NEXT: entry:
2741 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]]
2742 // FINITEONLY-NEXT: ret float [[CALL_I]]
2744 extern "C" __device__ float test_scalbnf(float x, int y) {
2745 return scalbnf(x, y);
2748 // DEFAULT-LABEL: @test_scalbn(
2749 // DEFAULT-NEXT: entry:
2750 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]]
2751 // DEFAULT-NEXT: ret double [[CALL_I]]
2753 // FINITEONLY-LABEL: @test_scalbn(
2754 // FINITEONLY-NEXT: entry:
2755 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR13]]
2756 // FINITEONLY-NEXT: ret double [[CALL_I]]
2758 extern "C" __device__ double test_scalbn(double x, int y) {
2759 return scalbn(x, y);
2762 // CHECK-LABEL: @test___signbitf(
2763 // CHECK-NEXT: entry:
2764 // CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f32(float noundef [[X:%.*]]) #[[ATTR13]]
2765 // CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
2766 // CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
2767 // CHECK-NEXT: ret i32 [[CONV]]
2769 extern "C" __device__ BOOL_TYPE test___signbitf(float x) {
2770 return __signbitf(x);
2773 // CHECK-LABEL: @test___signbit(
2774 // CHECK-NEXT: entry:
2775 // CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f64(double noundef [[X:%.*]]) #[[ATTR13]]
2776 // CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
2777 // CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
2778 // CHECK-NEXT: ret i32 [[CONV]]
2780 extern "C" __device__ BOOL_TYPE test___signbit(double x) {
2781 return __signbit(x);
2784 // DEFAULT-LABEL: @test_sincosf(
2785 // DEFAULT-NEXT: entry:
2786 // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
2787 // DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2788 // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
2789 // DEFAULT-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]]
2790 // DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15]]
2791 // DEFAULT-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]]
2792 // DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2793 // DEFAULT-NEXT: ret void
2795 // FINITEONLY-LABEL: @test_sincosf(
2796 // FINITEONLY-NEXT: entry:
2797 // FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
2798 // FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2799 // FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
2800 // FINITEONLY-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]]
2801 // FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15]]
2802 // FINITEONLY-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]]
2803 // FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2804 // FINITEONLY-NEXT: ret void
2806 extern "C" __device__ void test_sincosf(float x, float *y, float *z) {
2810 // DEFAULT-LABEL: @test_sincos(
2811 // DEFAULT-NEXT: entry:
2812 // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
2813 // DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2814 // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
2815 // DEFAULT-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]]
2816 // DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17]]
2817 // DEFAULT-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA17]]
2818 // DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2819 // DEFAULT-NEXT: ret void
2821 // FINITEONLY-LABEL: @test_sincos(
2822 // FINITEONLY-NEXT: entry:
2823 // FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
2824 // FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2825 // FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
2826 // FINITEONLY-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]]
2827 // FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17]]
2828 // FINITEONLY-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA17]]
2829 // FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2830 // FINITEONLY-NEXT: ret void
2832 extern "C" __device__ void test_sincos(double x, double *y, double *z) {
2836 // DEFAULT-LABEL: @test_sincospif(
2837 // DEFAULT-NEXT: entry:
2838 // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
2839 // DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2840 // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
2841 // DEFAULT-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]]
2842 // DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15]]
2843 // DEFAULT-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]]
2844 // DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2845 // DEFAULT-NEXT: ret void
2847 // FINITEONLY-LABEL: @test_sincospif(
2848 // FINITEONLY-NEXT: entry:
2849 // FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
2850 // FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2851 // FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
2852 // FINITEONLY-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]]
2853 // FINITEONLY-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA15]]
2854 // FINITEONLY-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]]
2855 // FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2856 // FINITEONLY-NEXT: ret void
2858 extern "C" __device__ void test_sincospif(float x, float *y, float *z) {
2862 // DEFAULT-LABEL: @test_sincospi(
2863 // DEFAULT-NEXT: entry:
2864 // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
2865 // DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2866 // DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
2867 // DEFAULT-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]]
2868 // DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17]]
2869 // DEFAULT-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA17]]
2870 // DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2871 // DEFAULT-NEXT: ret void
2873 // FINITEONLY-LABEL: @test_sincospi(
2874 // FINITEONLY-NEXT: entry:
2875 // FINITEONLY-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
2876 // FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2877 // FINITEONLY-NEXT: [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR15]]
2878 // FINITEONLY-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA17]]
2879 // FINITEONLY-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA17]]
2880 // FINITEONLY-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA17]]
2881 // FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR16]]
2882 // FINITEONLY-NEXT: ret void
2884 extern "C" __device__ void test_sincospi(double x, double *y, double *z) {
2888 // DEFAULT-LABEL: @test_sinf(
2889 // DEFAULT-NEXT: entry:
2890 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR15]]
2891 // DEFAULT-NEXT: ret float [[CALL_I]]
2893 // FINITEONLY-LABEL: @test_sinf(
2894 // FINITEONLY-NEXT: entry:
2895 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR15]]
2896 // FINITEONLY-NEXT: ret float [[CALL_I]]
2898 extern "C" __device__ float test_sinf(float x) {
2902 // DEFAULT-LABEL: @test_sin(
2903 // DEFAULT-NEXT: entry:
2904 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sin_f64(double noundef [[X:%.*]]) #[[ATTR15]]
2905 // DEFAULT-NEXT: ret double [[CALL_I]]
2907 // FINITEONLY-LABEL: @test_sin(
2908 // FINITEONLY-NEXT: entry:
2909 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sin_f64(double noundef [[X:%.*]]) #[[ATTR15]]
2910 // FINITEONLY-NEXT: ret double [[CALL_I]]
2912 extern "C" __device__ double test_sin(double x) {
2916 // DEFAULT-LABEL: @test_sinpif(
2917 // DEFAULT-NEXT: entry:
2918 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sinpi_f32(float noundef [[X:%.*]]) #[[ATTR15]]
2919 // DEFAULT-NEXT: ret float [[CALL_I]]
2921 // FINITEONLY-LABEL: @test_sinpif(
2922 // FINITEONLY-NEXT: entry:
2923 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sinpi_f32(float noundef [[X:%.*]]) #[[ATTR15]]
2924 // FINITEONLY-NEXT: ret float [[CALL_I]]
2926 extern "C" __device__ float test_sinpif(float x) {
2930 // DEFAULT-LABEL: @test_sinpi(
2931 // DEFAULT-NEXT: entry:
2932 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sinpi_f64(double noundef [[X:%.*]]) #[[ATTR15]]
2933 // DEFAULT-NEXT: ret double [[CALL_I]]
2935 // FINITEONLY-LABEL: @test_sinpi(
2936 // FINITEONLY-NEXT: entry:
2937 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sinpi_f64(double noundef [[X:%.*]]) #[[ATTR15]]
2938 // FINITEONLY-NEXT: ret double [[CALL_I]]
2940 extern "C" __device__ double test_sinpi(double x) {
2944 // DEFAULT-LABEL: @test_sqrtf(
2945 // DEFAULT-NEXT: entry:
2946 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR13]]
2947 // DEFAULT-NEXT: ret float [[CALL_I]]
2949 // FINITEONLY-LABEL: @test_sqrtf(
2950 // FINITEONLY-NEXT: entry:
2951 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR13]]
2952 // FINITEONLY-NEXT: ret float [[CALL_I]]
2954 extern "C" __device__ float test_sqrtf(float x) {
2958 // DEFAULT-LABEL: @test_sqrt(
2959 // DEFAULT-NEXT: entry:
2960 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR13]]
2961 // DEFAULT-NEXT: ret double [[CALL_I]]
2963 // FINITEONLY-LABEL: @test_sqrt(
2964 // FINITEONLY-NEXT: entry:
2965 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR13]]
2966 // FINITEONLY-NEXT: ret double [[CALL_I]]
2968 extern "C" __device__ double test_sqrt(double x) {
2972 // DEFAULT-LABEL: @test_tanf(
2973 // DEFAULT-NEXT: entry:
2974 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR15]]
2975 // DEFAULT-NEXT: ret float [[CALL_I]]
2977 // FINITEONLY-LABEL: @test_tanf(
2978 // FINITEONLY-NEXT: entry:
2979 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR15]]
2980 // FINITEONLY-NEXT: ret float [[CALL_I]]
2982 extern "C" __device__ float test_tanf(float x) {
2986 // DEFAULT-LABEL: @test_tan(
2987 // DEFAULT-NEXT: entry:
2988 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_tan_f64(double noundef [[X:%.*]]) #[[ATTR15]]
2989 // DEFAULT-NEXT: ret double [[CALL_I]]
2991 // FINITEONLY-LABEL: @test_tan(
2992 // FINITEONLY-NEXT: entry:
2993 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_tan_f64(double noundef [[X:%.*]]) #[[ATTR15]]
2994 // FINITEONLY-NEXT: ret double [[CALL_I]]
2996 extern "C" __device__ double test_tan(double x) {
3000 // DEFAULT-LABEL: @test_tanhf(
3001 // DEFAULT-NEXT: entry:
3002 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tanh_f32(float noundef [[X:%.*]]) #[[ATTR14]]
3003 // DEFAULT-NEXT: ret float [[CALL_I]]
3005 // FINITEONLY-LABEL: @test_tanhf(
3006 // FINITEONLY-NEXT: entry:
3007 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tanh_f32(float noundef [[X:%.*]]) #[[ATTR14]]
3008 // FINITEONLY-NEXT: ret float [[CALL_I]]
3010 extern "C" __device__ float test_tanhf(float x) {
3014 // DEFAULT-LABEL: @test_tanh(
3015 // DEFAULT-NEXT: entry:
3016 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_tanh_f64(double noundef [[X:%.*]]) #[[ATTR14]]
3017 // DEFAULT-NEXT: ret double [[CALL_I]]
3019 // FINITEONLY-LABEL: @test_tanh(
3020 // FINITEONLY-NEXT: entry:
3021 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_tanh_f64(double noundef [[X:%.*]]) #[[ATTR14]]
3022 // FINITEONLY-NEXT: ret double [[CALL_I]]
3024 extern "C" __device__ double test_tanh(double x) {
3028 // DEFAULT-LABEL: @test_tgammaf(
3029 // DEFAULT-NEXT: entry:
3030 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tgamma_f32(float noundef [[X:%.*]]) #[[ATTR15]]
3031 // DEFAULT-NEXT: ret float [[CALL_I]]
3033 // FINITEONLY-LABEL: @test_tgammaf(
3034 // FINITEONLY-NEXT: entry:
3035 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tgamma_f32(float noundef [[X:%.*]]) #[[ATTR15]]
3036 // FINITEONLY-NEXT: ret float [[CALL_I]]
3038 extern "C" __device__ float test_tgammaf(float x) {
3042 // DEFAULT-LABEL: @test_tgamma(
3043 // DEFAULT-NEXT: entry:
3044 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_tgamma_f64(double noundef [[X:%.*]]) #[[ATTR15]]
3045 // DEFAULT-NEXT: ret double [[CALL_I]]
3047 // FINITEONLY-LABEL: @test_tgamma(
3048 // FINITEONLY-NEXT: entry:
3049 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_tgamma_f64(double noundef [[X:%.*]]) #[[ATTR15]]
3050 // FINITEONLY-NEXT: ret double [[CALL_I]]
3052 extern "C" __device__ double test_tgamma(double x) {
3056 // DEFAULT-LABEL: @test_truncf(
3057 // DEFAULT-NEXT: entry:
3058 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_trunc_f32(float noundef [[X:%.*]]) #[[ATTR13]]
3059 // DEFAULT-NEXT: ret float [[CALL_I]]
3061 // FINITEONLY-LABEL: @test_truncf(
3062 // FINITEONLY-NEXT: entry:
3063 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_trunc_f32(float noundef [[X:%.*]]) #[[ATTR13]]
3064 // FINITEONLY-NEXT: ret float [[CALL_I]]
3066 extern "C" __device__ float test_truncf(float x) {
3070 // DEFAULT-LABEL: @test_trunc(
3071 // DEFAULT-NEXT: entry:
3072 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_trunc_f64(double noundef [[X:%.*]]) #[[ATTR13]]
3073 // DEFAULT-NEXT: ret double [[CALL_I]]
3075 // FINITEONLY-LABEL: @test_trunc(
3076 // FINITEONLY-NEXT: entry:
3077 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_trunc_f64(double noundef [[X:%.*]]) #[[ATTR13]]
3078 // FINITEONLY-NEXT: ret double [[CALL_I]]
3080 extern "C" __device__ double test_trunc(double x) {
3084 // DEFAULT-LABEL: @test_y0f(
3085 // DEFAULT-NEXT: entry:
3086 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[X:%.*]]) #[[ATTR15]]
3087 // DEFAULT-NEXT: ret float [[CALL_I]]
3089 // FINITEONLY-LABEL: @test_y0f(
3090 // FINITEONLY-NEXT: entry:
3091 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[X:%.*]]) #[[ATTR15]]
3092 // FINITEONLY-NEXT: ret float [[CALL_I]]
3094 extern "C" __device__ float test_y0f(float x) {
3098 // DEFAULT-LABEL: @test_y0(
3099 // DEFAULT-NEXT: entry:
3100 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[X:%.*]]) #[[ATTR15]]
3101 // DEFAULT-NEXT: ret double [[CALL_I]]
3103 // FINITEONLY-LABEL: @test_y0(
3104 // FINITEONLY-NEXT: entry:
3105 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[X:%.*]]) #[[ATTR15]]
3106 // FINITEONLY-NEXT: ret double [[CALL_I]]
3108 extern "C" __device__ double test_y0(double x) {
3112 // DEFAULT-LABEL: @test_y1f(
3113 // DEFAULT-NEXT: entry:
3114 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[X:%.*]]) #[[ATTR15]]
3115 // DEFAULT-NEXT: ret float [[CALL_I]]
3117 // FINITEONLY-LABEL: @test_y1f(
3118 // FINITEONLY-NEXT: entry:
3119 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[X:%.*]]) #[[ATTR15]]
3120 // FINITEONLY-NEXT: ret float [[CALL_I]]
3122 extern "C" __device__ float test_y1f(float x) {
3126 // DEFAULT-LABEL: @test_y1(
3127 // DEFAULT-NEXT: entry:
3128 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[X:%.*]]) #[[ATTR15]]
3129 // DEFAULT-NEXT: ret double [[CALL_I]]
3131 // FINITEONLY-LABEL: @test_y1(
3132 // FINITEONLY-NEXT: entry:
3133 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[X:%.*]]) #[[ATTR15]]
3134 // FINITEONLY-NEXT: ret double [[CALL_I]]
3136 extern "C" __device__ double test_y1(double x) {
3140 // DEFAULT-LABEL: @test_ynf(
3141 // DEFAULT-NEXT: entry:
3142 // DEFAULT-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
3143 // DEFAULT-NEXT: i32 0, label [[IF_THEN_I:%.*]]
3144 // DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
3146 // DEFAULT: if.then.i:
3147 // DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
3148 // DEFAULT-NEXT: br label [[_ZL3YNFIF_EXIT:%.*]]
3149 // DEFAULT: if.then2.i:
3150 // DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
3151 // DEFAULT-NEXT: br label [[_ZL3YNFIF_EXIT]]
3152 // DEFAULT: if.end4.i:
3153 // DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR15]]
3154 // DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
3155 // DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
3156 // DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
3157 // DEFAULT: for.body.i:
3158 // DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
3159 // DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
3160 // DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
3161 // DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
3162 // DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
3163 // DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
3164 // DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]]
3165 // DEFAULT-NEXT: [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]]
3166 // DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
3167 // DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
3168 // DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
3169 // DEFAULT: _ZL3ynfif.exit:
3170 // DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
3171 // DEFAULT-NEXT: ret float [[RETVAL_0_I]]
3173 // FINITEONLY-LABEL: @test_ynf(
3174 // FINITEONLY-NEXT: entry:
3175 // FINITEONLY-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
3176 // FINITEONLY-NEXT: i32 0, label [[IF_THEN_I:%.*]]
3177 // FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
3178 // FINITEONLY-NEXT: ]
3179 // FINITEONLY: if.then.i:
3180 // FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR15]]
3181 // FINITEONLY-NEXT: br label [[_ZL3YNFIF_EXIT:%.*]]
3182 // FINITEONLY: if.then2.i:
3183 // FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
3184 // FINITEONLY-NEXT: br label [[_ZL3YNFIF_EXIT]]
3185 // FINITEONLY: if.end4.i:
3186 // FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR15]]
3187 // FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR15]]
3188 // FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
3189 // FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
3190 // FINITEONLY: for.body.i:
3191 // FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
3192 // FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
3193 // FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
3194 // FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
3195 // FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
3196 // FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract float [[CONV_I]], [[Y]]
3197 // FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_025_I]], [[DIV_I]]
3198 // FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_024_I]]
3199 // FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
3200 // FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
3201 // FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
3202 // FINITEONLY: _ZL3ynfif.exit:
3203 // FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi float [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
3204 // FINITEONLY-NEXT: ret float [[RETVAL_0_I]]
3206 extern "C" __device__ float test_ynf(int x, float y) {
3210 // DEFAULT-LABEL: @test_yn(
3211 // DEFAULT-NEXT: entry:
3212 // DEFAULT-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
3213 // DEFAULT-NEXT: i32 0, label [[IF_THEN_I:%.*]]
3214 // DEFAULT-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
3216 // DEFAULT: if.then.i:
3217 // DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
3218 // DEFAULT-NEXT: br label [[_ZL2YNID_EXIT:%.*]]
3219 // DEFAULT: if.then2.i:
3220 // DEFAULT-NEXT: [[CALL_I20_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
3221 // DEFAULT-NEXT: br label [[_ZL2YNID_EXIT]]
3222 // DEFAULT: if.end4.i:
3223 // DEFAULT-NEXT: [[CALL_I21_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR15]]
3224 // DEFAULT-NEXT: [[CALL_I22_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
3225 // DEFAULT-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
3226 // DEFAULT-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
3227 // DEFAULT: for.body.i:
3228 // DEFAULT-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
3229 // DEFAULT-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
3230 // DEFAULT-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
3231 // DEFAULT-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
3232 // DEFAULT-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
3233 // DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
3234 // DEFAULT-NEXT: [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]]
3235 // DEFAULT-NEXT: [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]]
3236 // DEFAULT-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
3237 // DEFAULT-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
3238 // DEFAULT-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
3239 // DEFAULT: _ZL2ynid.exit:
3240 // DEFAULT-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
3241 // DEFAULT-NEXT: ret double [[RETVAL_0_I]]
3243 // FINITEONLY-LABEL: @test_yn(
3244 // FINITEONLY-NEXT: entry:
3245 // FINITEONLY-NEXT: switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
3246 // FINITEONLY-NEXT: i32 0, label [[IF_THEN_I:%.*]]
3247 // FINITEONLY-NEXT: i32 1, label [[IF_THEN2_I:%.*]]
3248 // FINITEONLY-NEXT: ]
3249 // FINITEONLY: if.then.i:
3250 // FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR15]]
3251 // FINITEONLY-NEXT: br label [[_ZL2YNID_EXIT:%.*]]
3252 // FINITEONLY: if.then2.i:
3253 // FINITEONLY-NEXT: [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
3254 // FINITEONLY-NEXT: br label [[_ZL2YNID_EXIT]]
3255 // FINITEONLY: if.end4.i:
3256 // FINITEONLY-NEXT: [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR15]]
3257 // FINITEONLY-NEXT: [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR15]]
3258 // FINITEONLY-NEXT: [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
3259 // FINITEONLY-NEXT: br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
3260 // FINITEONLY: for.body.i:
3261 // FINITEONLY-NEXT: [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
3262 // FINITEONLY-NEXT: [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
3263 // FINITEONLY-NEXT: [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
3264 // FINITEONLY-NEXT: [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
3265 // FINITEONLY-NEXT: [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
3266 // FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract double [[CONV_I]], [[Y]]
3267 // FINITEONLY-NEXT: [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_025_I]], [[DIV_I]]
3268 // FINITEONLY-NEXT: [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_024_I]]
3269 // FINITEONLY-NEXT: [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
3270 // FINITEONLY-NEXT: [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
3271 // FINITEONLY-NEXT: br i1 [[EXITCOND_NOT_I]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
3272 // FINITEONLY: _ZL2ynid.exit:
3273 // FINITEONLY-NEXT: [[RETVAL_0_I:%.*]] = phi double [ [[CALL_I_I]], [[IF_THEN_I]] ], [ [[CALL_I20_I]], [[IF_THEN2_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ], [ [[SUB_I]], [[FOR_BODY_I]] ]
3274 // FINITEONLY-NEXT: ret double [[RETVAL_0_I]]
3276 extern "C" __device__ double test_yn(int x, double y) {
3280 // DEFAULT-LABEL: @test___cosf(
3281 // DEFAULT-NEXT: entry:
3282 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR15]]
3283 // DEFAULT-NEXT: ret float [[CALL_I]]
3285 // FINITEONLY-LABEL: @test___cosf(
3286 // FINITEONLY-NEXT: entry:
3287 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR15]]
3288 // FINITEONLY-NEXT: ret float [[CALL_I]]
3290 extern "C" __device__ float test___cosf(float x) {
3294 // DEFAULT-LABEL: @test___exp10f(
3295 // DEFAULT-NEXT: entry:
3296 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_exp10_f32(float noundef [[X:%.*]]) #[[ATTR14]]
3297 // DEFAULT-NEXT: ret float [[CALL_I]]
3299 // FINITEONLY-LABEL: @test___exp10f(
3300 // FINITEONLY-NEXT: entry:
3301 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_exp10_f32(float noundef [[X:%.*]]) #[[ATTR14]]
3302 // FINITEONLY-NEXT: ret float [[CALL_I]]
3304 extern "C" __device__ float test___exp10f(float x) {
3308 // DEFAULT-LABEL: @test___expf(
3309 // DEFAULT-NEXT: entry:
3310 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_exp_f32(float noundef [[X:%.*]]) #[[ATTR14]]
3311 // DEFAULT-NEXT: ret float [[CALL_I]]
3313 // FINITEONLY-LABEL: @test___expf(
3314 // FINITEONLY-NEXT: entry:
3315 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_exp_f32(float noundef [[X:%.*]]) #[[ATTR14]]
3316 // FINITEONLY-NEXT: ret float [[CALL_I]]
3318 extern "C" __device__ float test___expf(float x) {
3322 // DEFAULT-LABEL: @test___fadd_rn(
3323 // DEFAULT-NEXT: entry:
3324 // DEFAULT-NEXT: [[ADD_I:%.*]] = fadd contract float [[X:%.*]], [[Y:%.*]]
3325 // DEFAULT-NEXT: ret float [[ADD_I]]
3327 // FINITEONLY-LABEL: @test___fadd_rn(
3328 // FINITEONLY-NEXT: entry:
3329 // FINITEONLY-NEXT: [[ADD_I:%.*]] = fadd nnan ninf contract float [[X:%.*]], [[Y:%.*]]
3330 // FINITEONLY-NEXT: ret float [[ADD_I]]
3332 extern "C" __device__ float test___fadd_rn(float x, float y) {
3333 return __fadd_rn(x, y);
3336 // DEFAULT-LABEL: @test___fdividef(
3337 // DEFAULT-NEXT: entry:
3338 // DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float [[X:%.*]], [[Y:%.*]]
3339 // DEFAULT-NEXT: ret float [[DIV_I]]
3341 // FINITEONLY-LABEL: @test___fdividef(
3342 // FINITEONLY-NEXT: entry:
3343 // FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract float [[X:%.*]], [[Y:%.*]]
3344 // FINITEONLY-NEXT: ret float [[DIV_I]]
3346 extern "C" __device__ float test___fdividef(float x, float y) {
3347 return __fdividef(x, y);
3350 // DEFAULT-LABEL: @test__fmaf_rn(
3351 // DEFAULT-NEXT: entry:
3352 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR13]]
3353 // DEFAULT-NEXT: ret float [[CALL_I]]
3355 // FINITEONLY-LABEL: @test__fmaf_rn(
3356 // FINITEONLY-NEXT: entry:
3357 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR13]]
3358 // FINITEONLY-NEXT: ret float [[CALL_I]]
3360 extern "C" __device__ float test__fmaf_rn(float x, float y, float z) {
3361 return __fmaf_rn(x, y, z);
3364 // DEFAULT-LABEL: @test___fmul_rn(
3365 // DEFAULT-NEXT: entry:
3366 // DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract float [[X:%.*]], [[Y:%.*]]
3367 // DEFAULT-NEXT: ret float [[MUL_I]]
3369 // FINITEONLY-LABEL: @test___fmul_rn(
3370 // FINITEONLY-NEXT: entry:
3371 // FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[X:%.*]], [[Y:%.*]]
3372 // FINITEONLY-NEXT: ret float [[MUL_I]]
3374 extern "C" __device__ float test___fmul_rn(float x, float y) {
3375 return __fmul_rn(x, y);
3378 // DEFAULT-LABEL: @test___frcp_rn(
3379 // DEFAULT-NEXT: entry:
3380 // DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract float 1.000000e+00, [[X:%.*]]
3381 // DEFAULT-NEXT: ret float [[DIV_I]]
3383 // FINITEONLY-LABEL: @test___frcp_rn(
3384 // FINITEONLY-NEXT: entry:
3385 // FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract float 1.000000e+00, [[X:%.*]]
3386 // FINITEONLY-NEXT: ret float [[DIV_I]]
3388 extern "C" __device__ float test___frcp_rn(float x) {
3389 return __frcp_rn(x);
3392 // DEFAULT-LABEL: @test___frsqrt_rn(
3393 // DEFAULT-NEXT: entry:
3394 // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.rsq.f32(float [[X:%.*]])
3395 // DEFAULT-NEXT: ret float [[TMP0]]
3397 // FINITEONLY-LABEL: @test___frsqrt_rn(
3398 // FINITEONLY-NEXT: entry:
3399 // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.amdgcn.rsq.f32(float [[X:%.*]])
3400 // FINITEONLY-NEXT: ret float [[TMP0]]
3402 extern "C" __device__ float test___frsqrt_rn(float x) {
3403 return __frsqrt_rn(x);
3406 // DEFAULT-LABEL: @test___fsqrt_rn(
3407 // DEFAULT-NEXT: entry:
3408 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR13]]
3409 // DEFAULT-NEXT: ret float [[CALL_I]]
3411 // FINITEONLY-LABEL: @test___fsqrt_rn(
3412 // FINITEONLY-NEXT: entry:
3413 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR13]]
3414 // FINITEONLY-NEXT: ret float [[CALL_I]]
3416 extern "C" __device__ float test___fsqrt_rn(float x) {
3417 return __fsqrt_rn(x);
3420 // DEFAULT-LABEL: @test___fsub_rn(
3421 // DEFAULT-NEXT: entry:
3422 // DEFAULT-NEXT: [[SUB_I:%.*]] = fsub contract float [[X:%.*]], [[Y:%.*]]
3423 // DEFAULT-NEXT: ret float [[SUB_I]]
3425 // FINITEONLY-LABEL: @test___fsub_rn(
3426 // FINITEONLY-NEXT: entry:
3427 // FINITEONLY-NEXT: [[SUB_I:%.*]] = fsub nnan ninf contract float [[X:%.*]], [[Y:%.*]]
3428 // FINITEONLY-NEXT: ret float [[SUB_I]]
3430 extern "C" __device__ float test___fsub_rn(float x, float y) {
3431 return __fsub_rn(x, y);
3434 // DEFAULT-LABEL: @test___log10f(
3435 // DEFAULT-NEXT: entry:
3436 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_log10_f32(float noundef [[X:%.*]]) #[[ATTR14]]
3437 // DEFAULT-NEXT: ret float [[CALL_I]]
3439 // FINITEONLY-LABEL: @test___log10f(
3440 // FINITEONLY-NEXT: entry:
3441 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_log10_f32(float noundef [[X:%.*]]) #[[ATTR14]]
3442 // FINITEONLY-NEXT: ret float [[CALL_I]]
3444 extern "C" __device__ float test___log10f(float x) {
3448 // DEFAULT-LABEL: @test___log2f(
3449 // DEFAULT-NEXT: entry:
3450 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_log2_f32(float noundef [[X:%.*]]) #[[ATTR14]]
3451 // DEFAULT-NEXT: ret float [[CALL_I]]
3453 // FINITEONLY-LABEL: @test___log2f(
3454 // FINITEONLY-NEXT: entry:
3455 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_log2_f32(float noundef [[X:%.*]]) #[[ATTR14]]
3456 // FINITEONLY-NEXT: ret float [[CALL_I]]
3458 extern "C" __device__ float test___log2f(float x) {
3462 // DEFAULT-LABEL: @test___logf(
3463 // DEFAULT-NEXT: entry:
3464 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_log_f32(float noundef [[X:%.*]]) #[[ATTR14]]
3465 // DEFAULT-NEXT: ret float [[CALL_I]]
3467 // FINITEONLY-LABEL: @test___logf(
3468 // FINITEONLY-NEXT: entry:
3469 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_log_f32(float noundef [[X:%.*]]) #[[ATTR14]]
3470 // FINITEONLY-NEXT: ret float [[CALL_I]]
3472 extern "C" __device__ float test___logf(float x) {
3476 // DEFAULT-LABEL: @test___powf(
3477 // DEFAULT-NEXT: entry:
3478 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
3479 // DEFAULT-NEXT: ret float [[CALL_I]]
3481 // FINITEONLY-LABEL: @test___powf(
3482 // FINITEONLY-NEXT: entry:
3483 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
3484 // FINITEONLY-NEXT: ret float [[CALL_I]]
3486 extern "C" __device__ float test___powf(float x, float y) {
3487 return __powf(x, y);
3490 // DEFAULT-LABEL: @test___saturatef(
3491 // DEFAULT-NEXT: entry:
3492 // DEFAULT-NEXT: [[CMP_I:%.*]] = fcmp contract olt float [[X:%.*]], 0.000000e+00
3493 // DEFAULT-NEXT: [[CMP1_I:%.*]] = fcmp contract ogt float [[X]], 1.000000e+00
3494 // DEFAULT-NEXT: [[COND_I:%.*]] = select contract i1 [[CMP1_I]], float 1.000000e+00, float [[X]]
3495 // DEFAULT-NEXT: [[COND5_I:%.*]] = select contract i1 [[CMP_I]], float 0.000000e+00, float [[COND_I]]
3496 // DEFAULT-NEXT: ret float [[COND5_I]]
3498 // FINITEONLY-LABEL: @test___saturatef(
3499 // FINITEONLY-NEXT: entry:
3500 // FINITEONLY-NEXT: [[CMP_I:%.*]] = fcmp nnan ninf contract olt float [[X:%.*]], 0.000000e+00
3501 // FINITEONLY-NEXT: [[CMP1_I:%.*]] = fcmp nnan ninf contract ogt float [[X]], 1.000000e+00
3502 // FINITEONLY-NEXT: [[COND_I:%.*]] = select nnan ninf contract i1 [[CMP1_I]], float 1.000000e+00, float [[X]]
3503 // FINITEONLY-NEXT: [[COND5_I:%.*]] = select nnan ninf contract i1 [[CMP_I]], float 0.000000e+00, float [[COND_I]]
3504 // FINITEONLY-NEXT: ret float [[COND5_I]]
3506 extern "C" __device__ float test___saturatef(float x) {
3507 return __saturatef(x);
3510 // DEFAULT-LABEL: @test___sincosf(
3511 // DEFAULT-NEXT: entry:
3512 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR15]]
3513 // DEFAULT-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]]
3514 // DEFAULT-NEXT: [[CALL1_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR15]]
3515 // DEFAULT-NEXT: store float [[CALL1_I]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]]
3516 // DEFAULT-NEXT: ret void
3518 // FINITEONLY-LABEL: @test___sincosf(
3519 // FINITEONLY-NEXT: entry:
3520 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR15]]
3521 // FINITEONLY-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA15]]
3522 // FINITEONLY-NEXT: [[CALL1_I:%.*]] = tail call nnan ninf contract float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR15]]
3523 // FINITEONLY-NEXT: store float [[CALL1_I]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA15]]
3524 // FINITEONLY-NEXT: ret void
3526 extern "C" __device__ void test___sincosf(float x, float *y, float *z) {
3530 // DEFAULT-LABEL: @test___sinf(
3531 // DEFAULT-NEXT: entry:
3532 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR15]]
3533 // DEFAULT-NEXT: ret float [[CALL_I]]
3535 // FINITEONLY-LABEL: @test___sinf(
3536 // FINITEONLY-NEXT: entry:
3537 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR15]]
3538 // FINITEONLY-NEXT: ret float [[CALL_I]]
3540 extern "C" __device__ float test___sinf(float x) {
3544 // DEFAULT-LABEL: @test___tanf(
3545 // DEFAULT-NEXT: entry:
3546 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR15]]
3547 // DEFAULT-NEXT: ret float [[CALL_I]]
3549 // FINITEONLY-LABEL: @test___tanf(
3550 // FINITEONLY-NEXT: entry:
3551 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR15]]
3552 // FINITEONLY-NEXT: ret float [[CALL_I]]
3554 extern "C" __device__ float test___tanf(float x) {
3558 // DEFAULT-LABEL: @test___dadd_rn(
3559 // DEFAULT-NEXT: entry:
3560 // DEFAULT-NEXT: [[ADD_I:%.*]] = fadd contract double [[X:%.*]], [[Y:%.*]]
3561 // DEFAULT-NEXT: ret double [[ADD_I]]
3563 // FINITEONLY-LABEL: @test___dadd_rn(
3564 // FINITEONLY-NEXT: entry:
3565 // FINITEONLY-NEXT: [[ADD_I:%.*]] = fadd nnan ninf contract double [[X:%.*]], [[Y:%.*]]
3566 // FINITEONLY-NEXT: ret double [[ADD_I]]
3568 extern "C" __device__ double test___dadd_rn(double x, double y) {
3569 return __dadd_rn(x, y);
3572 // DEFAULT-LABEL: @test___ddiv_rn(
3573 // DEFAULT-NEXT: entry:
3574 // DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract double [[X:%.*]], [[Y:%.*]]
3575 // DEFAULT-NEXT: ret double [[DIV_I]]
3577 // FINITEONLY-LABEL: @test___ddiv_rn(
3578 // FINITEONLY-NEXT: entry:
3579 // FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract double [[X:%.*]], [[Y:%.*]]
3580 // FINITEONLY-NEXT: ret double [[DIV_I]]
3582 extern "C" __device__ double test___ddiv_rn(double x, double y) {
3583 return __ddiv_rn(x, y);
3586 // DEFAULT-LABEL: @test___dmul_rn(
3587 // DEFAULT-NEXT: entry:
3588 // DEFAULT-NEXT: [[MUL_I:%.*]] = fmul contract double [[X:%.*]], [[Y:%.*]]
3589 // DEFAULT-NEXT: ret double [[MUL_I]]
3591 // FINITEONLY-LABEL: @test___dmul_rn(
3592 // FINITEONLY-NEXT: entry:
3593 // FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract double [[X:%.*]], [[Y:%.*]]
3594 // FINITEONLY-NEXT: ret double [[MUL_I]]
3596 extern "C" __device__ double test___dmul_rn(double x, double y) {
3597 return __dmul_rn(x, y);
3600 // DEFAULT-LABEL: @test___drcp_rn(
3601 // DEFAULT-NEXT: entry:
3602 // DEFAULT-NEXT: [[DIV_I:%.*]] = fdiv contract double 1.000000e+00, [[X:%.*]]
3603 // DEFAULT-NEXT: ret double [[DIV_I]]
3605 // FINITEONLY-LABEL: @test___drcp_rn(
3606 // FINITEONLY-NEXT: entry:
3607 // FINITEONLY-NEXT: [[DIV_I:%.*]] = fdiv nnan ninf contract double 1.000000e+00, [[X:%.*]]
3608 // FINITEONLY-NEXT: ret double [[DIV_I]]
3610 extern "C" __device__ double test___drcp_rn(double x) {
3611 return __drcp_rn(x);
3614 // DEFAULT-LABEL: @test___dsqrt_rn(
3615 // DEFAULT-NEXT: entry:
3616 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR13]]
3617 // DEFAULT-NEXT: ret double [[CALL_I]]
3619 // FINITEONLY-LABEL: @test___dsqrt_rn(
3620 // FINITEONLY-NEXT: entry:
3621 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR13]]
3622 // FINITEONLY-NEXT: ret double [[CALL_I]]
3624 extern "C" __device__ double test___dsqrt_rn(double x) {
3625 return __dsqrt_rn(x);
3628 // DEFAULT-LABEL: @test__fma_rn(
3629 // DEFAULT-NEXT: entry:
3630 // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR13]]
3631 // DEFAULT-NEXT: ret double [[CALL_I]]
3633 // FINITEONLY-LABEL: @test__fma_rn(
3634 // FINITEONLY-NEXT: entry:
3635 // FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR13]]
3636 // FINITEONLY-NEXT: ret double [[CALL_I]]
3638 extern "C" __device__ double test__fma_rn(double x, double y, double z) {
3639 return __fma_rn(x, y, z);
3642 // DEFAULT-LABEL: @test_float_min(
3643 // DEFAULT-NEXT: entry:
3644 // DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
3645 // DEFAULT-NEXT: ret float [[CALL_I_I]]
3647 // FINITEONLY-LABEL: @test_float_min(
3648 // FINITEONLY-NEXT: entry:
3649 // FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
3650 // FINITEONLY-NEXT: ret float [[CALL_I_I]]
3652 extern "C" __device__ float test_float_min(float x, float y) {
3656 // DEFAULT-LABEL: @test_float_max(
3657 // DEFAULT-NEXT: entry:
3658 // DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
3659 // DEFAULT-NEXT: ret float [[CALL_I_I]]
3661 // FINITEONLY-LABEL: @test_float_max(
3662 // FINITEONLY-NEXT: entry:
3663 // FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR13]]
3664 // FINITEONLY-NEXT: ret float [[CALL_I_I]]
3666 extern "C" __device__ float test_float_max(float x, float y) {
3670 // DEFAULT-LABEL: @test_double_min(
3671 // DEFAULT-NEXT: entry:
3672 // DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
3673 // DEFAULT-NEXT: ret double [[CALL_I_I]]
3675 // FINITEONLY-LABEL: @test_double_min(
3676 // FINITEONLY-NEXT: entry:
3677 // FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
3678 // FINITEONLY-NEXT: ret double [[CALL_I_I]]
3680 extern "C" __device__ double test_double_min(double x, double y) {
3684 // DEFAULT-LABEL: @test_double_max(
3685 // DEFAULT-NEXT: entry:
3686 // DEFAULT-NEXT: [[CALL_I_I:%.*]] = tail call contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
3687 // DEFAULT-NEXT: ret double [[CALL_I_I]]
3689 // FINITEONLY-LABEL: @test_double_max(
3690 // FINITEONLY-NEXT: entry:
3691 // FINITEONLY-NEXT: [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR13]]
3692 // FINITEONLY-NEXT: ret double [[CALL_I_I]]
3694 extern "C" __device__ double test_double_max(double x, double y) {
3697 // CHECK-LABEL: @test_int_min(
3698 // CHECK-NEXT: entry:
3699 // CHECK-NEXT: [[COND_I:%.*]] = tail call i32 @llvm.smin.i32(i32 [[X:%.*]], i32 [[Y:%.*]])
3700 // CHECK-NEXT: ret i32 [[COND_I]]
3702 extern "C" __device__ int test_int_min(int x, int y) {
3706 // CHECK-LABEL: @test_int_max(
3707 // CHECK-NEXT: entry:
3708 // CHECK-NEXT: [[COND_I:%.*]] = tail call i32 @llvm.smax.i32(i32 [[X:%.*]], i32 [[Y:%.*]])
3709 // CHECK-NEXT: ret i32 [[COND_I]]
3711 extern "C" __device__ int test_int_max(int x, int y) {