clang/HIP: Add new header test for math IR gen

[arsenm]

Summary

The current header testing is pretty thin. This is in
preparation for a series of patches to replace many
builtin implementations.

I did try to stress everything in this header, but skipped
a few things. Mostly I didn't understand why we have
various language version checks which skip defining some
things. It doesn't seem right to have any of these if guards
on cplusplus, HIPCC_RTC, and OPENMP_AMDGCN__.

[arsenm]

Missed frexp

[JonChesterfield]

Good call adding this test, thanks.

I think some of the skipping is glibc quirk related and some was to avoid clobbering the user's namespace for openmp. The math header defines various functions at global scope which can collide with user symbols. In practice that may not be a big deal.

[yaxunl]

LGTM. Thanks.

[arsenm]

2edafe8393fa7c9ad2e95d691191469adb7bf4b0

clang/test/Headers/__clang_hip_math.hip

+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
+// REQUIRES: amdgpu-registered-target
+
+// Test without OCML_BASIC_ROUNDED_OPERATIONS
+// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
+// RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
+// RUN:   -internal-isystem %S/Inputs/include \
+// RUN:   -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
+// RUN:   -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -o - \
+// RUN:   -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,DEFAULT %s
+
+// Check that we end up with fast math flags set on intrinsic calls
+// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
+// RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
+// RUN:   -internal-isystem %S/Inputs/include \
+// RUN:   -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
+// RUN:   -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -ffinite-math-only -o - \
+// RUN:   -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,FINITEONLY %s
+
+#define BOOL_TYPE int
+typedef unsigned long long uint64_t;
+
+// CHECK-LABEL: @test___make_mantissa_base8(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    br label [[WHILE_COND_I:%.*]]
+// CHECK:       while.cond.i:
+// CHECK-NEXT:    [[__TAGP_ADDR_0_I:%.*]] = phi ptr [ [[P:%.*]], [[ENTRY:%.*]] ], [ [[__TAGP_ADDR_1_I:%.*]], [[CLEANUP_I:%.*]] ]
+// CHECK-NEXT:    [[__R_0_I:%.*]] = phi i64 [ 0, [[ENTRY]] ], [ [[__R_1_I:%.*]], [[CLEANUP_I]] ]
+// CHECK-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq ptr [[__TAGP_ADDR_0_I]], null
+// CHECK-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL21__MAKE_MANTISSA_BASE8PKC_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// CHECK:       while.body.i:
+// CHECK-NEXT:    [[TMP0:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I]], align 1, !tbaa [[TBAA3:![0-9]+]]
+// CHECK-NEXT:    [[TMP1:%.*]] = and i8 [[TMP0]], -8
+// CHECK-NEXT:    [[TMP2:%.*]] = icmp eq i8 [[TMP1]], 48
+// CHECK-NEXT:    br i1 [[TMP2]], label [[IF_THEN_I:%.*]], label [[CLEANUP_I]]
+// CHECK:       if.then.i:
+// CHECK-NEXT:    [[MUL_I:%.*]] = shl i64 [[__R_0_I]], 3
+// CHECK-NEXT:    [[CONV3_I:%.*]] = sext i8 [[TMP0]] to i64
+// CHECK-NEXT:    [[ADD_I:%.*]] = add i64 [[MUL_I]], -48
+// CHECK-NEXT:    [[SUB_I:%.*]] = add i64 [[ADD_I]], [[CONV3_I]]
+// CHECK-NEXT:    [[INCDEC_PTR_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I]], i64 1
+// CHECK-NEXT:    br label [[CLEANUP_I]]
+// CHECK:       cleanup.i:
+// CHECK-NEXT:    [[__TAGP_ADDR_1_I]] = phi ptr [ [[INCDEC_PTR_I]], [[IF_THEN_I]] ], [ [[__TAGP_ADDR_0_I]], [[WHILE_BODY_I]] ]
+// CHECK-NEXT:    [[__R_1_I]] = phi i64 [ [[SUB_I]], [[IF_THEN_I]] ], [ [[__R_0_I]], [[WHILE_BODY_I]] ]
+// CHECK-NEXT:    br i1 [[TMP2]], label [[WHILE_COND_I]], label [[_ZL21__MAKE_MANTISSA_BASE8PKC_EXIT]], !llvm.loop [[LOOP6:![0-9]+]]
+// CHECK:       _ZL21__make_mantissa_base8PKc.exit:
+// CHECK-NEXT:    [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ]
+// CHECK-NEXT:    ret i64 [[RETVAL_2_I]]
+//
+extern "C" __device__ uint64_t test___make_mantissa_base8(const char *p) {
+  return __make_mantissa_base8(p);
+}
+
+// CHECK-LABEL: @test___make_mantissa_base10(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    br label [[WHILE_COND_I:%.*]]
+// CHECK:       while.cond.i:
+// CHECK-NEXT:    [[__TAGP_ADDR_0_I:%.*]] = phi ptr [ [[P:%.*]], [[ENTRY:%.*]] ], [ [[__TAGP_ADDR_1_I:%.*]], [[CLEANUP_I:%.*]] ]
+// CHECK-NEXT:    [[__R_0_I:%.*]] = phi i64 [ 0, [[ENTRY]] ], [ [[__R_1_I:%.*]], [[CLEANUP_I]] ]
+// CHECK-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq ptr [[__TAGP_ADDR_0_I]], null
+// CHECK-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL22__MAKE_MANTISSA_BASE10PKC_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// CHECK:       while.body.i:
+// CHECK-NEXT:    [[TMP0:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT:    [[TMP1:%.*]] = add i8 [[TMP0]], -48
+// CHECK-NEXT:    [[TMP2:%.*]] = icmp ult i8 [[TMP1]], 10
+// CHECK-NEXT:    br i1 [[TMP2]], label [[IF_THEN_I:%.*]], label [[CLEANUP_I]]
+// CHECK:       if.then.i:
+// CHECK-NEXT:    [[MUL_I:%.*]] = mul i64 [[__R_0_I]], 10
+// CHECK-NEXT:    [[CONV3_I:%.*]] = sext i8 [[TMP0]] to i64
+// CHECK-NEXT:    [[ADD_I:%.*]] = add i64 [[MUL_I]], -48
+// CHECK-NEXT:    [[SUB_I:%.*]] = add i64 [[ADD_I]], [[CONV3_I]]
+// CHECK-NEXT:    [[INCDEC_PTR_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I]], i64 1
+// CHECK-NEXT:    br label [[CLEANUP_I]]
+// CHECK:       cleanup.i:
+// CHECK-NEXT:    [[__TAGP_ADDR_1_I]] = phi ptr [ [[INCDEC_PTR_I]], [[IF_THEN_I]] ], [ [[__TAGP_ADDR_0_I]], [[WHILE_BODY_I]] ]
+// CHECK-NEXT:    [[__R_1_I]] = phi i64 [ [[SUB_I]], [[IF_THEN_I]] ], [ [[__R_0_I]], [[WHILE_BODY_I]] ]
+// CHECK-NEXT:    br i1 [[TMP2]], label [[WHILE_COND_I]], label [[_ZL22__MAKE_MANTISSA_BASE10PKC_EXIT]], !llvm.loop [[LOOP9:![0-9]+]]
+// CHECK:       _ZL22__make_mantissa_base10PKc.exit:
+// CHECK-NEXT:    [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ]
+// CHECK-NEXT:    ret i64 [[RETVAL_2_I]]
+//
+extern "C" __device__ uint64_t test___make_mantissa_base10(const char *p) {
+  return __make_mantissa_base10(p);
+}
+
+// CHECK-LABEL: @test___make_mantissa_base16(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    br label [[WHILE_COND_I:%.*]]
+// CHECK:       while.cond.i:
+// CHECK-NEXT:    [[__TAGP_ADDR_0_I:%.*]] = phi ptr [ [[P:%.*]], [[ENTRY:%.*]] ], [ [[__TAGP_ADDR_1_I:%.*]], [[CLEANUP_I:%.*]] ]
+// CHECK-NEXT:    [[__R_0_I:%.*]] = phi i64 [ 0, [[ENTRY]] ], [ [[__R_2_I:%.*]], [[CLEANUP_I]] ]
+// CHECK-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq ptr [[__TAGP_ADDR_0_I]], null
+// CHECK-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// CHECK:       while.body.i:
+// CHECK-NEXT:    [[TMP0:%.*]] = load i8, ptr [[__TAGP_ADDR_0_I]], align 1, !tbaa [[TBAA3]]
+// CHECK-NEXT:    [[TMP1:%.*]] = add i8 [[TMP0]], -48
+// CHECK-NEXT:    [[TMP2:%.*]] = icmp ult i8 [[TMP1]], 10
+// CHECK-NEXT:    br i1 [[TMP2]], label [[IF_END29_I:%.*]], label [[IF_ELSE_I:%.*]]
+// CHECK:       if.else.i:
+// CHECK-NEXT:    [[TMP3:%.*]] = add i8 [[TMP0]], -97
+// CHECK-NEXT:    [[TMP4:%.*]] = icmp ult i8 [[TMP3]], 6
+// CHECK-NEXT:    br i1 [[TMP4]], label [[IF_END29_I]], label [[IF_ELSE15_I:%.*]]
+// CHECK:       if.else15.i:
+// CHECK-NEXT:    [[TMP5:%.*]] = add i8 [[TMP0]], -65
+// CHECK-NEXT:    [[TMP6:%.*]] = icmp ult i8 [[TMP5]], 6
+// CHECK-NEXT:    br i1 [[TMP6]], label [[IF_END29_I]], label [[CLEANUP_I]]
+// CHECK:       if.end29.i:
+// CHECK-NEXT:    [[DOTSINK:%.*]] = phi i64 [ -48, [[WHILE_BODY_I]] ], [ -87, [[IF_ELSE_I]] ], [ -55, [[IF_ELSE15_I]] ]
+// CHECK-NEXT:    [[MUL22_I:%.*]] = shl i64 [[__R_0_I]], 4
+// CHECK-NEXT:    [[CONV23_I:%.*]] = sext i8 [[TMP0]] to i64
+// CHECK-NEXT:    [[ADD24_I:%.*]] = add i64 [[MUL22_I]], [[DOTSINK]]
+// CHECK-NEXT:    [[ADD26_I:%.*]] = add i64 [[ADD24_I]], [[CONV23_I]]
+// CHECK-NEXT:    [[INCDEC_PTR_I:%.*]] = getelementptr inbounds i8, ptr [[__TAGP_ADDR_0_I]], i64 1
+// CHECK-NEXT:    br label [[CLEANUP_I]]
+// CHECK:       cleanup.i:
+// CHECK-NEXT:    [[__TAGP_ADDR_1_I]] = phi ptr [ [[INCDEC_PTR_I]], [[IF_END29_I]] ], [ [[__TAGP_ADDR_0_I]], [[IF_ELSE15_I]] ]
+// CHECK-NEXT:    [[__R_2_I]] = phi i64 [ [[ADD26_I]], [[IF_END29_I]] ], [ [[__R_0_I]], [[IF_ELSE15_I]] ]
+// CHECK-NEXT:    [[COND_I:%.*]] = phi i1 [ true, [[IF_END29_I]] ], [ false, [[IF_ELSE15_I]] ]
+// CHECK-NEXT:    br i1 [[COND_I]], label [[WHILE_COND_I]], label [[_ZL22__MAKE_MANTISSA_BASE16PKC_EXIT]], !llvm.loop [[LOOP10:![0-9]+]]
+// CHECK:       _ZL22__make_mantissa_base16PKc.exit:
+// CHECK-NEXT:    [[RETVAL_2_I:%.*]] = phi i64 [ 0, [[CLEANUP_I]] ], [ [[__R_0_I]], [[WHILE_COND_I]] ]
+// CHECK-NEXT:    ret i64 [[RETVAL_2_I]]
+//
+extern "C" __device__ uint64_t test___make_mantissa_base16(const char *p) {
+  return __make_mantissa_base16(p);
+}
+
+// CHECK-LABEL: @test___make_mantissa(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret i64 0
+//
+extern "C" __device__ uint64_t test___make_mantissa(const char *p) {
+  return __make_mantissa(p);
+}
+
+// CHECK-LABEL: @test_abs(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[TMP0:%.*]] = tail call i32 @llvm.abs.i32(i32 [[X:%.*]], i1 true)
+// CHECK-NEXT:    ret i32 [[TMP0]]
+//
+extern "C" __device__ int test_abs(int x) {
+  return abs(x);
+}
+
+// CHECK-LABEL: @test_labs(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[TMP0:%.*]] = tail call i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true)
+// CHECK-NEXT:    ret i64 [[TMP0]]
+//
+extern "C" __device__ long test_labs(long x) {
+  return labs(x);
+}
+
+//
+// CHECK-LABEL: @test_llabs(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[TMP0:%.*]] = tail call i64 @llvm.abs.i64(i64 [[X:%.*]], i1 true)
+// CHECK-NEXT:    ret i64 [[TMP0]]
+//
+extern "C" __device__ long long test_llabs(long x) {
+  return llabs(x);
+}
+
+// DEFAULT-LABEL: @test_acosf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_acos_f32(float noundef [[X:%.*]]) #[[ATTR14:[0-9]+]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_acosf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_acos_f32(float noundef [[X:%.*]]) #[[ATTR14:[0-9]+]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_acosf(float x) {
+  return acosf(x);
+}
+
+// DEFAULT-LABEL: @test_acos(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_acos_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_acos(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_acos_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_acos(double x) {
+  return acos(x);
+}
+
+// DEFAULT-LABEL: @test_acoshf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_acosh_f32(float noundef [[X:%.*]]) #[[ATTR15:[0-9]+]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_acoshf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_acosh_f32(float noundef [[X:%.*]]) #[[ATTR15:[0-9]+]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_acoshf(float x) {
+  return acoshf(x);
+}
+
+// DEFAULT-LABEL: @test_acosh(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_acosh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_acosh(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_acosh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_acosh(double x) {
+  return acosh(x);
+}
+
+// DEFAULT-LABEL: @test_asinf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_asin_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_asinf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_asin_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_asinf(float x) {
+  return asinf(x);
+}
+
+// DEFAULT-LABEL: @test_asin(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_asin_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_asin(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_asin_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_asin(double x) {
+
+  return asin(x);
+}
+
+// DEFAULT-LABEL: @test_asinhf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_asinh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_asinhf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_asinh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_asinhf(float x) {
+  return asinhf(x);
+}
+
+// DEFAULT-LABEL: @test_asinh(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_asinh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_asinh(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_asinh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_asinh(double x) {
+  return asinh(x);
+}
+
+// DEFAULT-LABEL: @test_atan2f(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_atan2_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_atan2f(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_atan2_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_atan2f(float x, float y) {
+  return atan2f(x, y);
+}
+
+// DEFAULT-LABEL: @test_atan2(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_atan2_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_atan2(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_atan2_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_atan2(double x, double y) {
+  return atan2(x, y);
+}
+
+// DEFAULT-LABEL: @test_atanf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_atan_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_atanf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_atan_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_atanf(float x) {
+  return atanf(x);
+}
+
+// DEFAULT-LABEL: @test_atan(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_atan_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_atan(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_atan_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_atan(double x) {
+  return atan(x);
+}
+
+// DEFAULT-LABEL: @test_atanhf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_atanh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_atanhf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_atanh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_atanhf(float x) {
+  return atanhf(x);
+}
+
+// DEFAULT-LABEL: @test_atanh(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_atanh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_atanh(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_atanh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_atanh(double x) {
+  return atanh(x);
+}
+
+// DEFAULT-LABEL: @test_cbrtf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_cbrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_cbrtf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cbrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_cbrtf(float x) {
+  return cbrtf(x);
+}
+
+// DEFAULT-LABEL: @test_cbrt(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_cbrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_cbrt(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cbrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_cbrt(double x) {
+  return cbrt(x);
+}
+
+// DEFAULT-LABEL: @test_ceilf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_ceil_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_ceilf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ceil_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_ceilf(float x) {
+  return ceilf(x);
+}
+
+// DEFAULT-LABEL: @test_ceil(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_ceil_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_ceil(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ceil_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_ceil(double x) {
+  return ceil(x);
+}
+
+// DEFAULT-LABEL: @test_copysignf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_copysign_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_copysignf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_copysign_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_copysignf(float x, float y) {
+  return copysignf(x, y);
+}
+
+// DEFAULT-LABEL: @test_copysign(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_copysign_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_copysign(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_copysign_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_copysign(double x, double y) {
+  return copysign(x, y);
+}
+
+// DEFAULT-LABEL: @test_cosf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR16:[0-9]+]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_cosf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR16:[0-9]+]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_cosf(float x) {
+  return cosf(x);
+}
+
+// DEFAULT-LABEL: @test_cos(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_cos_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_cos(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cos_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_cos(double x) {
+  return cos(x);
+}
+
+// DEFAULT-LABEL: @test_coshf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_cosh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_coshf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cosh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_coshf(float x) {
+  return coshf(x);
+}
+
+// DEFAULT-LABEL: @test_cosh(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_cosh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_cosh(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cosh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_cosh(double x) {
+  return cosh(x);
+}
+
+// DEFAULT-LABEL: @test_cospif(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_cospi_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_cospif(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_cospi_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_cospif(float x) {
+  return cospif(x);
+}
+
+// DEFAULT-LABEL: @test_cospi(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_cospi_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_cospi(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_cospi_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_cospi(double x) {
+  return cospi(x);
+}
+
+// DEFAULT-LABEL: @test_cyl_bessel_i0f(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_i0_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_cyl_bessel_i0f(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_i0_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_cyl_bessel_i0f(float x) {
+  return cyl_bessel_i0f(x);
+}
+
+// DEFAULT-LABEL: @test_cyl_bessel_i0(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_i0_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_cyl_bessel_i0(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_i0_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_cyl_bessel_i0(double x) {
+  return cyl_bessel_i0(x);
+}
+
+// DEFAULT-LABEL: @test_cyl_bessel_i1f(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_i1_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_cyl_bessel_i1f(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_i1_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_cyl_bessel_i1f(float x) {
+  return cyl_bessel_i1f(x);
+}
+
+// DEFAULT-LABEL: @test_cyl_bessel_i1(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_i1_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_cyl_bessel_i1(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_i1_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_cyl_bessel_i1(double x) {
+  return cyl_bessel_i1(x);
+}
+
+// DEFAULT-LABEL: @test_erfcf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_erfc_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_erfcf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_erfc_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_erfcf(float x) {
+  return erfcf(x);
+}
+
+// DEFAULT-LABEL: @test_erfc(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_erfc_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_erfc(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_erfc_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_erfc(double x) {
+  return erfc(x);
+}
+
+// DEFAULT-LABEL: @test_erfinvf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_erfinv_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_erfinvf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_erfinv_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_erfinvf(float x) {
+  return erfinvf(x);
+}
+
+// DEFAULT-LABEL: @test_erfinv(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_erfinv_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_erfinv(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_erfinv_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_erfinv(double x) {
+  return erfinv(x);
+}
+
+// DEFAULT-LABEL: @test_exp10f(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_exp10_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_exp10f(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_exp10_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_exp10f(float x) {
+  return exp10f(x);
+}
+
+// DEFAULT-LABEL: @test_exp10(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_exp10_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_exp10(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_exp10_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_exp10(double x) {
+  return exp10(x);
+}
+
+// DEFAULT-LABEL: @test_exp2f(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_exp2_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_exp2f(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_exp2_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_exp2f(float x) {
+  return exp2f(x);
+}
+
+// DEFAULT-LABEL: @test_exp2(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_exp2(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_exp2(double x) {
+  return exp2(x);
+}
+
+// DEFAULT-LABEL: @test_expf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_exp_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_expf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_exp_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_expf(float x) {
+  return expf(x);
+}
+
+// DEFAULT-LABEL: @test_exp(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_exp_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_exp(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_exp_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_exp(double x) {
+  return exp(x);
+}
+
+// DEFAULT-LABEL: @test_expm1f(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_expm1_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_expm1f(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_expm1_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_expm1f(float x) {
+  return expm1f(x);
+}
+
+// DEFAULT-LABEL: @test_expm1(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_expm1_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_expm1(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_expm1_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_expm1(double x) {
+  return expm1(x);
+}
+
+// DEFAULT-LABEL: @test_fabsf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_fabs_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_fabsf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fabs_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_fabsf(float x) {
+  return fabsf(x);
+}
+
+// DEFAULT-LABEL: @test_fabs(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_fabs_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_fabs(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fabs_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_fabs(double x) {
+  return fabs(x);
+}
+
+// DEFAULT-LABEL: @test_fdimf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_fdim_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_fdimf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fdim_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_fdimf(float x, float y) {
+  return fdimf(x, y);
+}
+
+// DEFAULT-LABEL: @test_fdim(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_fdim_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_fdim(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fdim_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_fdim(double x, double y) {
+  return fdim(x, y);
+}
+
+// DEFAULT-LABEL: @test_fdividef(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[DIV_I:%.*]] = fdiv contract float [[X:%.*]], [[Y:%.*]]
+// DEFAULT-NEXT:    ret float [[DIV_I]]
+//
+// FINITEONLY-LABEL: @test_fdividef(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[DIV_I:%.*]] = fdiv nnan ninf contract float [[X:%.*]], [[Y:%.*]]
+// FINITEONLY-NEXT:    ret float [[DIV_I]]
+//
+extern "C" __device__ float test_fdividef(float x, float y) {
+  return fdividef(x, y);
+}
+
+// DEFAULT-LABEL: @test_floorf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_floor_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_floorf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_floor_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_floorf(float x) {
+  return floorf(x);
+}
+
+// DEFAULT-LABEL: @test_floor(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_floor_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_floor(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_floor_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_floor(double x) {
+  return floor(x);
+}
+
+// DEFAULT-LABEL: @test_fmaf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_fmaf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_fmaf(float x, float y, float z) {
+  return fmaf(x, y, z);
+}
+
+// DEFAULT-LABEL: @test_fma(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_fma(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_fma(double x, double y, double z) {
+  return fma(x, y, z);
+}
+
+// DEFAULT-LABEL: @test_fma_rn(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_fma_rn(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_fma_rn(double x, double y, double z) {
+  return __fma_rn(x, y, z);
+}
+
+// DEFAULT-LABEL: @test_fmaxf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_fmaxf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_fmaxf(float x, float y) {
+  return fmaxf(x, y);
+}
+
+// DEFAULT-LABEL: @test_fmax(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_fmax(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_fmax(double x, double y) {
+  return fmax(x, y);
+}
+
+// DEFAULT-LABEL: @test_fminf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_fminf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_fminf(float x, float y) {
+  return fminf(x, y);
+}
+
+// DEFAULT-LABEL: @test_fmin(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_fmin(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_fmin(double x, double y) {
+  return fmin(x, y);
+}
+
+// DEFAULT-LABEL: @test_fmodf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_fmod_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_fmodf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fmod_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_fmodf(float x, float y) {
+  return fmodf(x, y);
+}
+
+// DEFAULT-LABEL: @test_fmod(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_fmod_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_fmod(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fmod_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_fmod(double x, double y) {
+  return fmod(x, y);
+}
+
+// DEFAULT-LABEL: @test_hypotf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_hypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_hypotf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_hypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_hypotf(float x, float y) {
+  return hypotf(x, y);
+}
+
+// DEFAULT-LABEL: @test_hypot(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_hypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_hypot(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_hypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_hypot(double x, double y) {
+  return hypot(x, y);
+}
+
+// CHECK-LABEL: @test_ilogbf(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f32(float noundef [[X:%.*]]) #[[ATTR14:[0-9]+]]
+// CHECK-NEXT:    ret i32 [[CALL_I]]
+//
+extern "C" __device__ int test_ilogbf(float x) {
+  return ilogbf(x);
+}
+
+// CHECK-LABEL: @test_ilogb(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// CHECK-NEXT:    ret i32 [[CALL_I]]
+//
+extern "C" __device__ int test_ilogb(double x) {
+  return ilogb(x);
+}
+
+// CHECK-LABEL: @test___finitef(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_isfinite_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// CHECK-NEXT:    [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
+// CHECK-NEXT:    [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
+// CHECK-NEXT:    ret i32 [[CONV]]
+//
+extern "C" __device__ BOOL_TYPE test___finitef(float x) {
+  return __finitef(x);
+}
+
+// CHECK-LABEL: @test___finite(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_isfinite_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// CHECK-NEXT:    [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
+// CHECK-NEXT:    [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
+// CHECK-NEXT:    ret i32 [[CONV]]
+//
+extern "C" __device__ BOOL_TYPE test___finite(double x) {
+  return __finite(x);
+}
+
+// CHECK-LABEL: @test___isinff(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// CHECK-NEXT:    [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
+// CHECK-NEXT:    [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
+// CHECK-NEXT:    ret i32 [[CONV]]
+//
+extern "C" __device__ BOOL_TYPE test___isinff(float x) {
+  return __isinff(x);
+}
+
+// CHECK-LABEL: @test___isinf(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_isinf_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// CHECK-NEXT:    [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
+// CHECK-NEXT:    [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
+// CHECK-NEXT:    ret i32 [[CONV]]
+//
+extern "C" __device__ BOOL_TYPE test___isinf(double x) {
+  return __isinf(x);
+}
+
+// CHECK-LABEL: @test___isnanf(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// CHECK-NEXT:    [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
+// CHECK-NEXT:    [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
+// CHECK-NEXT:    ret i32 [[CONV]]
+//
+extern "C" __device__ BOOL_TYPE test___isnanf(float x) {
+  return __isnanf(x);
+}
+
+// CHECK-LABEL: @test___isnan(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// CHECK-NEXT:    [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
+// CHECK-NEXT:    [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
+// CHECK-NEXT:    ret i32 [[CONV]]
+//
+extern "C" __device__ BOOL_TYPE test___isnan(double x) {
+  return __isnan(x);
+}
+
+// DEFAULT-LABEL: @test_j0f(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_j0f(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_j0f(float x) {
+  return j0f(x);
+}
+
+// DEFAULT-LABEL: @test_j0(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_j0(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_j0(double x) {
+  return j0(x);
+}
+
+// DEFAULT-LABEL: @test_j1f(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_j1f(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_j1f(float x) {
+  return j1f(x);
+}
+
+// DEFAULT-LABEL: @test_j1(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_j1(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_j1(double x) {
+  return j1(x);
+}
+
+// DEFAULT-LABEL: @test_jnf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
+// DEFAULT-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
+// DEFAULT-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
+// DEFAULT-NEXT:    ]
+// DEFAULT:       if.then.i:
+// DEFAULT-NEXT:    [[CALL_I_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    br label [[_ZL3JNFIF_EXIT:%.*]]
+// DEFAULT:       if.then2.i:
+// DEFAULT-NEXT:    [[CALL_I20_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT:    br label [[_ZL3JNFIF_EXIT]]
+// DEFAULT:       if.end4.i:
+// DEFAULT-NEXT:    [[CALL_I21_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT:    [[CALL_I22_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
+// DEFAULT:       for.body.i:
+// DEFAULT-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
+// DEFAULT-NEXT:    [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
+// DEFAULT-NEXT:    [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]]
+// DEFAULT-NEXT:    [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]]
+// DEFAULT-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// DEFAULT-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP11:![0-9]+]]
+// DEFAULT:       _ZL3jnfif.exit:
+// 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]] ]
+// DEFAULT-NEXT:    ret float [[RETVAL_0_I]]
+//
+// FINITEONLY-LABEL: @test_jnf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
+// FINITEONLY-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
+// FINITEONLY-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
+// FINITEONLY-NEXT:    ]
+// FINITEONLY:       if.then.i:
+// FINITEONLY-NEXT:    [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    br label [[_ZL3JNFIF_EXIT:%.*]]
+// FINITEONLY:       if.then2.i:
+// FINITEONLY-NEXT:    [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    br label [[_ZL3JNFIF_EXIT]]
+// FINITEONLY:       if.end4.i:
+// FINITEONLY-NEXT:    [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
+// FINITEONLY:       for.body.i:
+// FINITEONLY-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
+// FINITEONLY-NEXT:    [[DIV_I:%.*]] = fdiv nnan ninf contract float [[CONV_I]], [[Y]]
+// FINITEONLY-NEXT:    [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_025_I]], [[DIV_I]]
+// FINITEONLY-NEXT:    [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_024_I]]
+// FINITEONLY-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// FINITEONLY-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP11:![0-9]+]]
+// FINITEONLY:       _ZL3jnfif.exit:
+// 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]] ]
+// FINITEONLY-NEXT:    ret float [[RETVAL_0_I]]
+//
+extern "C" __device__ float test_jnf(int x, float y) {
+  return jnf(x, y);
+}
+
+// DEFAULT-LABEL: @test_jn(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
+// DEFAULT-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
+// DEFAULT-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
+// DEFAULT-NEXT:    ]
+// DEFAULT:       if.then.i:
+// DEFAULT-NEXT:    [[CALL_I_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    br label [[_ZL2JNID_EXIT:%.*]]
+// DEFAULT:       if.then2.i:
+// DEFAULT-NEXT:    [[CALL_I20_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT:    br label [[_ZL2JNID_EXIT]]
+// DEFAULT:       if.end4.i:
+// DEFAULT-NEXT:    [[CALL_I21_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT:    [[CALL_I22_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
+// DEFAULT:       for.body.i:
+// DEFAULT-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
+// DEFAULT-NEXT:    [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
+// DEFAULT-NEXT:    [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]]
+// DEFAULT-NEXT:    [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]]
+// DEFAULT-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// DEFAULT-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP12:![0-9]+]]
+// DEFAULT:       _ZL2jnid.exit:
+// 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]] ]
+// DEFAULT-NEXT:    ret double [[RETVAL_0_I]]
+//
+// FINITEONLY-LABEL: @test_jn(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
+// FINITEONLY-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
+// FINITEONLY-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
+// FINITEONLY-NEXT:    ]
+// FINITEONLY:       if.then.i:
+// FINITEONLY-NEXT:    [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    br label [[_ZL2JNID_EXIT:%.*]]
+// FINITEONLY:       if.then2.i:
+// FINITEONLY-NEXT:    [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    br label [[_ZL2JNID_EXIT]]
+// FINITEONLY:       if.end4.i:
+// FINITEONLY-NEXT:    [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
+// FINITEONLY:       for.body.i:
+// FINITEONLY-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
+// FINITEONLY-NEXT:    [[DIV_I:%.*]] = fdiv nnan ninf contract double [[CONV_I]], [[Y]]
+// FINITEONLY-NEXT:    [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_025_I]], [[DIV_I]]
+// FINITEONLY-NEXT:    [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_024_I]]
+// FINITEONLY-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// FINITEONLY-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP12:![0-9]+]]
+// FINITEONLY:       _ZL2jnid.exit:
+// 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]] ]
+// FINITEONLY-NEXT:    ret double [[RETVAL_0_I]]
+//
+extern "C" __device__ double test_jn(int x, double y) {
+  return jn(x, y);
+}
+
+// DEFAULT-LABEL: @test_ldexpf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_ldexp_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_ldexpf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ldexp_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_ldexpf(float x, int y) {
+  return ldexpf(x, y);
+}
+
+// DEFAULT-LABEL: @test_ldexp(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_ldexp_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_ldexp(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ldexp_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_ldexp(double x, int y) {
+  return ldexp(x, y);
+}
+
+// DEFAULT-LABEL: @test_lgammaf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_lgamma_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_lgammaf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_lgamma_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_lgammaf(float x) {
+  return lgammaf(x);
+}
+
+// DEFAULT-LABEL: @test_lgamma(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_lgamma_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_lgamma(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_lgamma_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_lgamma(double x) {
+  return lgamma(x);
+}
+
+// DEFAULT-LABEL: @test_llrintf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
+// DEFAULT-NEXT:    ret i64 [[CONV_I]]
+//
+// FINITEONLY-LABEL: @test_llrintf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
+// FINITEONLY-NEXT:    ret i64 [[CONV_I]]
+//
+extern "C" __device__ long long int test_llrintf(float x) {
+  return llrintf(x);
+}
+
+// DEFAULT-LABEL: @test_llrint(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
+// DEFAULT-NEXT:    ret i64 [[CONV_I]]
+//
+// FINITEONLY-LABEL: @test_llrint(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
+// FINITEONLY-NEXT:    ret i64 [[CONV_I]]
+//
+extern "C" __device__ long long int test_llrint(double x) {
+  return llrint(x);
+}
+
+// DEFAULT-LABEL: @test_llroundf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
+// DEFAULT-NEXT:    ret i64 [[CONV_I]]
+//
+// FINITEONLY-LABEL: @test_llroundf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
+// FINITEONLY-NEXT:    ret i64 [[CONV_I]]
+//
+extern "C" __device__ long long int test_llroundf(float x) {
+  return llroundf(x);
+}
+
+// DEFAULT-LABEL: @test_llround(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
+// DEFAULT-NEXT:    ret i64 [[CONV_I]]
+//
+// FINITEONLY-LABEL: @test_llround(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
+// FINITEONLY-NEXT:    ret i64 [[CONV_I]]
+//
+extern "C" __device__ long long int test_llround(double x) {
+  return llround(x);
+}
+
+// DEFAULT-LABEL: @test_log10f(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_log10_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_log10f(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_log10_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_log10f(float x) {
+  return log10f(x);
+}
+
+// DEFAULT-LABEL: @test_log10(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_log10_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_log10(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_log10_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_log10(double x) {
+  return log10(x);
+}
+
+// DEFAULT-LABEL: @test_log1pf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_log1p_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_log1pf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_log1p_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_log1pf(float x) {
+  return log1pf(x);
+}
+
+// DEFAULT-LABEL: @test_log1p(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_log1p_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_log1p(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_log1p_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_log1p(double x) {
+  return log1p(x);
+}
+
+// DEFAULT-LABEL: @test_log2f(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_log2_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_log2f(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_log2_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_log2f(float x) {
+  return log2f(x);
+}
+
+// DEFAULT-LABEL: @test_log2(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_log2_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_log2(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_log2_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_log2(double x) {
+  return log2(x);
+}
+
+// DEFAULT-LABEL: @test_logbf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_logb_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_logbf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_logb_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_logbf(float x) {
+  return logbf(x);
+}
+
+// DEFAULT-LABEL: @test_logb(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_logb_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_logb(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_logb_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_logb(double x) {
+  return logb(x);
+}
+
+// DEFAULT-LABEL: @test_lrintf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
+// DEFAULT-NEXT:    ret i64 [[CONV_I]]
+//
+// FINITEONLY-LABEL: @test_lrintf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
+// FINITEONLY-NEXT:    ret i64 [[CONV_I]]
+//
+extern "C" __device__ long int test_lrintf(float x) {
+  return lrintf(x);
+}
+
+// DEFAULT-LABEL: @test_lrint(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
+// DEFAULT-NEXT:    ret i64 [[CONV_I]]
+//
+// FINITEONLY-LABEL: @test_lrint(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
+// FINITEONLY-NEXT:    ret i64 [[CONV_I]]
+//
+extern "C" __device__ long int test_lrint(double x) {
+  return lrint(x);
+}
+
+// DEFAULT-LABEL: @test_lroundf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
+// DEFAULT-NEXT:    ret i64 [[CONV_I]]
+//
+// FINITEONLY-LABEL: @test_lroundf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64
+// FINITEONLY-NEXT:    ret i64 [[CONV_I]]
+//
+extern "C" __device__ long int test_lroundf(float x) {
+  return lroundf(x);
+}
+
+// DEFAULT-LABEL: @test_lround(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
+// DEFAULT-NEXT:    ret i64 [[CONV_I]]
+//
+// FINITEONLY-LABEL: @test_lround(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64
+// FINITEONLY-NEXT:    ret i64 [[CONV_I]]
+//
+extern "C" __device__ long int test_lround(double x) {
+  return lround(x);
+}
+
+// DEFAULT-LABEL: @test_modff(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
+// DEFAULT-NEXT:    call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17:[0-9]+]]
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = call contract float @__ocml_modf_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// DEFAULT-NEXT:    [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13:![0-9]+]]
+// DEFAULT-NEXT:    store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]]
+// DEFAULT-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_modff(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
+// FINITEONLY-NEXT:    call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17:[0-9]+]]
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_modf_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13:![0-9]+]]
+// FINITEONLY-NEXT:    store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]]
+// FINITEONLY-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_modff(float x, float* y) {
+  return modff(x, y);
+}
+
+// DEFAULT-LABEL: @test_modf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
+// DEFAULT-NEXT:    call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = call contract double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// DEFAULT-NEXT:    [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15:![0-9]+]]
+// DEFAULT-NEXT:    store double [[TMP0]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]]
+// DEFAULT-NEXT:    call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_modf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
+// FINITEONLY-NEXT:    call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15:![0-9]+]]
+// FINITEONLY-NEXT:    store double [[TMP0]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]]
+// FINITEONLY-NEXT:    call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_modf(double x, double* y) {
+  return modf(x, y);
+}
+
+// CHECK-LABEL: @test_nanf(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret float 0x7FF8000000000000
+//
+extern "C" __device__ float test_nanf(const char *tag) {
+  return nanf(tag);
+}
+
+// CHECK-LABEL: @test_nan(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret double 0x7FF8000000000000
+//
+extern "C" __device__ double test_nan(const char *tag) {
+  return nan(tag);
+}
+
+// CHECK-LABEL: @test_nanf_emptystr(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret float 0x7FF8000000000000
+//
+extern "C" __device__ float test_nanf_emptystr() {
+  return nanf("");
+}
+
+// CHECK-LABEL: @test_nan_emptystr(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret double 0x7FF8000000000000
+//
+extern "C" __device__ double test_nan_emptystr() {
+  return nan("");
+}
+
+// CHECK-LABEL: @test_nanf_fill(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret float 0x7FF8000000000000
+//
+extern "C" __device__ float test_nanf_fill() {
+  return nanf("0x456");
+}
+
+// CHECK-LABEL: @test_nan_fill(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret double 0x7FF8000000000000
+//
+extern "C" __device__ double test_nan_fill() {
+  return nan("0x123");
+}
+
+// DEFAULT-LABEL: @test_nearbyintf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_nearbyint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_nearbyintf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_nearbyint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_nearbyintf(float x) {
+  return nearbyintf(x);
+}
+
+// DEFAULT-LABEL: @test_nearbyint(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_nearbyint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_nearbyint(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_nearbyint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_nearbyint(double x) {
+  return nearbyint(x);
+}
+
+// DEFAULT-LABEL: @test_nextafterf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_nextafter_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_nextafterf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_nextafter_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_nextafterf(float x, float y) {
+  return nextafterf(x, y);
+}
+
+// DEFAULT-LABEL: @test_nextafter(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_nextafter_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_nextafter(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_nextafter_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_nextafter(double x, double y) {
+  return nextafter(x, y);
+}
+
+// DEFAULT-LABEL: @test_norm3df(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_len3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_norm3df(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_len3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_norm3df(float x, float y, float z) {
+  return norm3df(x, y, z);
+}
+
+// DEFAULT-LABEL: @test_norm3d(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_len3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_norm3d(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_len3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_norm3d(double x, double y, double z) {
+  return norm3d(x, y, z);
+}
+
+// DEFAULT-LABEL: @test_norm4df(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_len4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_norm4df(
+// FINITEONLY-NEXT:  entry:
+// 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:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_norm4df(float x, float y, float z, float w) {
+  return norm4df(x, y, z, w);
+}
+
+// DEFAULT-LABEL: @test_norm4d(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_len4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_norm4d(
+// FINITEONLY-NEXT:  entry:
+// 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:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_norm4d(double x, double y, double z, double w) {
+  return norm4d(x, y, z, w);
+}
+
+// DEFAULT-LABEL: @test_normcdff(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_ncdf_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_normcdff(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ncdf_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_normcdff(float x) {
+  return normcdff(x);
+}
+
+// DEFAULT-LABEL: @test_normcdf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_ncdf_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_normcdf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ncdf_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_normcdf(double x) {
+  return normcdf(x);
+}
+
+// DEFAULT-LABEL: @test_normcdfinvf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_ncdfinv_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_normcdfinvf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_ncdfinv_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_normcdfinvf(float x) {
+  return normcdfinvf(x);
+}
+
+// DEFAULT-LABEL: @test_normcdfinv(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_ncdfinv_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_normcdfinv(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_ncdfinv_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_normcdfinv(double x) {
+  return normcdfinv(x);
+}
+
+// DEFAULT-LABEL: @test_normf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// DEFAULT:       while.body.i:
+// DEFAULT-NEXT:    [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// DEFAULT-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// DEFAULT-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// DEFAULT-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
+// DEFAULT-NEXT:    [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA13]]
+// DEFAULT-NEXT:    [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]]
+// DEFAULT-NEXT:    [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]]
+// DEFAULT-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// DEFAULT-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
+// DEFAULT-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP17:![0-9]+]]
+// DEFAULT:       _ZL5normfiPKf.exit:
+// DEFAULT-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_normf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// FINITEONLY:       while.body.i:
+// FINITEONLY-NEXT:    [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// FINITEONLY-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// FINITEONLY-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// FINITEONLY-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA13]]
+// FINITEONLY-NEXT:    [[MUL_I:%.*]] = fmul nnan ninf contract float [[TMP0]], [[TMP0]]
+// FINITEONLY-NEXT:    [[ADD_I]] = fadd nnan ninf contract float [[__R_08_I]], [[MUL_I]]
+// FINITEONLY-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// FINITEONLY-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
+// FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP17:![0-9]+]]
+// FINITEONLY:       _ZL5normfiPKf.exit:
+// FINITEONLY-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_normf(int x, const float *y) {
+  return normf(x, y);
+}
+
+// DEFAULT-LABEL: @test_norm(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// DEFAULT:       while.body.i:
+// DEFAULT-NEXT:    [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// DEFAULT-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// DEFAULT-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// DEFAULT-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
+// DEFAULT-NEXT:    [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA15]]
+// DEFAULT-NEXT:    [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]]
+// DEFAULT-NEXT:    [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]]
+// DEFAULT-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// DEFAULT-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
+// DEFAULT-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP18:![0-9]+]]
+// DEFAULT:       _ZL4normiPKd.exit:
+// DEFAULT-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_norm(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// FINITEONLY:       while.body.i:
+// FINITEONLY-NEXT:    [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// FINITEONLY-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// FINITEONLY-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// FINITEONLY-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA15]]
+// FINITEONLY-NEXT:    [[MUL_I:%.*]] = fmul nnan ninf contract double [[TMP0]], [[TMP0]]
+// FINITEONLY-NEXT:    [[ADD_I]] = fadd nnan ninf contract double [[__R_08_I]], [[MUL_I]]
+// FINITEONLY-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// FINITEONLY-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
+// FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP18:![0-9]+]]
+// FINITEONLY:       _ZL4normiPKd.exit:
+// FINITEONLY-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_norm(int x, const double *y) {
+  return norm(x, y);
+}
+
+// DEFAULT-LABEL: @test_powf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_powf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_powf(float x, float y) {
+  return powf(x, y);
+}
+
+// DEFAULT-LABEL: @test_pow(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_pow_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_pow(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_pow_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_pow(double x, double y) {
+  return pow(x, y);
+}
+
+// DEFAULT-LABEL: @test_powif(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_pown_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_powif(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_pown_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_powif(float x, int y) {
+  return powif(x, y);
+}
+
+// DEFAULT-LABEL: @test_powi(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_pown_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_powi(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_pown_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_powi(double x, int y) {
+  return powi(x, y);
+}
+
+// DEFAULT-LABEL: @test_rcbrtf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_rcbrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_rcbrtf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rcbrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_rcbrtf(float x) {
+  return rcbrtf(x);
+}
+
+// DEFAULT-LABEL: @test_rcbrt(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_rcbrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_rcbrt(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rcbrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_rcbrt(double x) {
+  return rcbrt(x);
+}
+
+// DEFAULT-LABEL: @test_remainderf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_remainder_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_remainderf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_remainder_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_remainderf(float x, float y) {
+  return remainderf(x, y);
+}
+
+// DEFAULT-LABEL: @test_remainder(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_remainder_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_remainder(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_remainder_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_remainder(double x, double y) {
+  return remainder(x, y);
+}
+
+// DEFAULT-LABEL: @test_remquof(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
+// DEFAULT-NEXT:    call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = call contract float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// DEFAULT-NEXT:    [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA19:![0-9]+]]
+// DEFAULT-NEXT:    store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA19]]
+// DEFAULT-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_remquof(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
+// FINITEONLY-NEXT:    call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA19:![0-9]+]]
+// FINITEONLY-NEXT:    store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA19]]
+// FINITEONLY-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_remquof(float x, float y, int* z) {
+  return remquof(x, y, z);
+}
+
+// DEFAULT-LABEL: @test_remquo(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
+// DEFAULT-NEXT:    call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = call contract double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// DEFAULT-NEXT:    [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA19]]
+// DEFAULT-NEXT:    store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA19]]
+// DEFAULT-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_remquo(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
+// FINITEONLY-NEXT:    call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA19]]
+// FINITEONLY-NEXT:    store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA19]]
+// FINITEONLY-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_remquo(double x, double y, int* z) {
+  return remquo(x, y, z);
+}
+
+// DEFAULT-LABEL: @test_rhypotf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_rhypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_rhypotf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rhypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_rhypotf(float x, float y) {
+  return rhypotf(x, y);
+}
+
+// DEFAULT-LABEL: @test_rhypot(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_rhypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_rhypot(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rhypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_rhypot(double x, double y) {
+  return rhypot(x, y);
+}
+
+// DEFAULT-LABEL: @test_rintf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_rintf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rint_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_rintf(float x) {
+  return rintf(x);
+}
+
+// DEFAULT-LABEL: @test_rint(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_rint(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rint_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_rint(double x) {
+  return rint(x);
+}
+
+// DEFAULT-LABEL: @test_rnormf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// DEFAULT:       while.body.i:
+// DEFAULT-NEXT:    [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// DEFAULT-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// DEFAULT-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// DEFAULT-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
+// DEFAULT-NEXT:    [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA13]]
+// DEFAULT-NEXT:    [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]]
+// DEFAULT-NEXT:    [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]]
+// DEFAULT-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// DEFAULT-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
+// DEFAULT-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
+// DEFAULT:       _ZL6rnormfiPKf.exit:
+// DEFAULT-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_rnormf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// FINITEONLY:       while.body.i:
+// FINITEONLY-NEXT:    [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// FINITEONLY-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// FINITEONLY-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// FINITEONLY-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA13]]
+// FINITEONLY-NEXT:    [[MUL_I:%.*]] = fmul nnan ninf contract float [[TMP0]], [[TMP0]]
+// FINITEONLY-NEXT:    [[ADD_I]] = fadd nnan ninf contract float [[__R_08_I]], [[MUL_I]]
+// FINITEONLY-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// FINITEONLY-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
+// FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
+// FINITEONLY:       _ZL6rnormfiPKf.exit:
+// FINITEONLY-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_rnormf(int x, const float* y) {
+  return rnormf(x, y);
+}
+
+// DEFAULT-LABEL: @test_rnorm(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// DEFAULT-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// DEFAULT:       while.body.i:
+// DEFAULT-NEXT:    [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// DEFAULT-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// DEFAULT-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// DEFAULT-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
+// DEFAULT-NEXT:    [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA15]]
+// DEFAULT-NEXT:    [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]]
+// DEFAULT-NEXT:    [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]]
+// DEFAULT-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// DEFAULT-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
+// DEFAULT-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]]
+// DEFAULT:       _ZL5rnormiPKd.exit:
+// DEFAULT-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_rnorm(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// FINITEONLY:       while.body.i:
+// FINITEONLY-NEXT:    [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// FINITEONLY-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// FINITEONLY-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// FINITEONLY-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA15]]
+// FINITEONLY-NEXT:    [[MUL_I:%.*]] = fmul nnan ninf contract double [[TMP0]], [[TMP0]]
+// FINITEONLY-NEXT:    [[ADD_I]] = fadd nnan ninf contract double [[__R_08_I]], [[MUL_I]]
+// FINITEONLY-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// FINITEONLY-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
+// FINITEONLY-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]]
+// FINITEONLY:       _ZL5rnormiPKd.exit:
+// FINITEONLY-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_rnorm(int x, const double* y) {
+  return rnorm(x, y);
+}
+
+// DEFAULT-LABEL: @test_rnorm3df(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_rlen3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_rnorm3df(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rlen3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_rnorm3df(float x, float y, float z) {
+  return rnorm3df(x, y, z);
+}
+
+// DEFAULT-LABEL: @test_rnorm3d(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_rlen3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_rnorm3d(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rlen3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_rnorm3d(double x, double y, double z) {
+  return rnorm3d(x, y, z);
+}
+
+// DEFAULT-LABEL: @test_rnorm4df(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_rlen4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_rnorm4df(
+// FINITEONLY-NEXT:  entry:
+// 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:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_rnorm4df(float x, float y, float z, float w) {
+  return rnorm4df(x, y, z, w);
+}
+
+// DEFAULT-LABEL: @test_rnorm4d(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_rlen4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_rnorm4d(
+// FINITEONLY-NEXT:  entry:
+// 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:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_rnorm4d(double x, double y, double z, double w) {
+  return rnorm4d(x, y, z, w);
+}
+
+// DEFAULT-LABEL: @test_roundf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_roundf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_roundf(float x) {
+  return roundf(x);
+}
+
+// DEFAULT-LABEL: @test_round(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_round(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_round(double x) {
+  return round(x);
+}
+
+// DEFAULT-LABEL: @test_rsqrtf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_rsqrtf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_rsqrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_rsqrtf(float x) {
+  return rsqrtf(x);
+}
+
+// DEFAULT-LABEL: @test_rsqrt(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_rsqrt(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_rsqrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_rsqrt(double x) {
+  return rsqrt(x);
+}
+
+// DEFAULT-LABEL: @test_scalblnf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CMP_NOT_I:%.*]] = icmp eq i64 [[Y:%.*]], 9223372036854775807
+// DEFAULT-NEXT:    br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]]
+// DEFAULT:       cond.true.i:
+// DEFAULT-NEXT:    [[CONV_I:%.*]] = trunc i64 [[Y]] to i32
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]]
+// DEFAULT-NEXT:    br label [[_ZL8SCALBLNFFL_EXIT:%.*]]
+// DEFAULT:       cond.false.i:
+// DEFAULT-NEXT:    [[CALL2_I:%.*]] = tail call contract float @__ocml_scalb_f32(float noundef [[X]], float noundef 0x43E0000000000000) #[[ATTR14]]
+// DEFAULT-NEXT:    br label [[_ZL8SCALBLNFFL_EXIT]]
+// DEFAULT:       _ZL8scalblnffl.exit:
+// DEFAULT-NEXT:    [[COND_I:%.*]] = phi contract float [ [[CALL_I]], [[COND_TRUE_I]] ], [ [[CALL2_I]], [[COND_FALSE_I]] ]
+// DEFAULT-NEXT:    ret float [[COND_I]]
+//
+// FINITEONLY-LABEL: @test_scalblnf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CMP_NOT_I:%.*]] = icmp eq i64 [[Y:%.*]], 9223372036854775807
+// FINITEONLY-NEXT:    br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]]
+// FINITEONLY:       cond.true.i:
+// FINITEONLY-NEXT:    [[CONV_I:%.*]] = trunc i64 [[Y]] to i32
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    br label [[_ZL8SCALBLNFFL_EXIT:%.*]]
+// FINITEONLY:       cond.false.i:
+// FINITEONLY-NEXT:    [[CALL2_I:%.*]] = tail call nnan ninf contract float @__ocml_scalb_f32(float noundef [[X]], float noundef 0x43E0000000000000) #[[ATTR14]]
+// FINITEONLY-NEXT:    br label [[_ZL8SCALBLNFFL_EXIT]]
+// FINITEONLY:       _ZL8scalblnffl.exit:
+// FINITEONLY-NEXT:    [[COND_I:%.*]] = phi nnan ninf contract float [ [[CALL_I]], [[COND_TRUE_I]] ], [ [[CALL2_I]], [[COND_FALSE_I]] ]
+// FINITEONLY-NEXT:    ret float [[COND_I]]
+//
+extern "C" __device__ float test_scalblnf(float x, long int y) {
+  return scalblnf(x, y);
+}
+
+// DEFAULT-LABEL: @test_scalbln(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CMP_NOT_I:%.*]] = icmp eq i64 [[Y:%.*]], 9223372036854775807
+// DEFAULT-NEXT:    br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]]
+// DEFAULT:       cond.true.i:
+// DEFAULT-NEXT:    [[CONV_I:%.*]] = trunc i64 [[Y]] to i32
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]]
+// DEFAULT-NEXT:    br label [[_ZL7SCALBLNDL_EXIT:%.*]]
+// DEFAULT:       cond.false.i:
+// DEFAULT-NEXT:    [[CALL2_I:%.*]] = tail call contract double @__ocml_scalb_f64(double noundef [[X]], double noundef 0x43E0000000000000) #[[ATTR14]]
+// DEFAULT-NEXT:    br label [[_ZL7SCALBLNDL_EXIT]]
+// DEFAULT:       _ZL7scalblndl.exit:
+// DEFAULT-NEXT:    [[COND_I:%.*]] = phi contract double [ [[CALL_I]], [[COND_TRUE_I]] ], [ [[CALL2_I]], [[COND_FALSE_I]] ]
+// DEFAULT-NEXT:    ret double [[COND_I]]
+//
+// FINITEONLY-LABEL: @test_scalbln(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CMP_NOT_I:%.*]] = icmp eq i64 [[Y:%.*]], 9223372036854775807
+// FINITEONLY-NEXT:    br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]]
+// FINITEONLY:       cond.true.i:
+// FINITEONLY-NEXT:    [[CONV_I:%.*]] = trunc i64 [[Y]] to i32
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[CONV_I]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    br label [[_ZL7SCALBLNDL_EXIT:%.*]]
+// FINITEONLY:       cond.false.i:
+// FINITEONLY-NEXT:    [[CALL2_I:%.*]] = tail call nnan ninf contract double @__ocml_scalb_f64(double noundef [[X]], double noundef 0x43E0000000000000) #[[ATTR14]]
+// FINITEONLY-NEXT:    br label [[_ZL7SCALBLNDL_EXIT]]
+// FINITEONLY:       _ZL7scalblndl.exit:
+// FINITEONLY-NEXT:    [[COND_I:%.*]] = phi nnan ninf contract double [ [[CALL_I]], [[COND_TRUE_I]] ], [ [[CALL2_I]], [[COND_FALSE_I]] ]
+// FINITEONLY-NEXT:    ret double [[COND_I]]
+//
+extern "C" __device__ double test_scalbln(double x, long int y) {
+  return scalbln(x, y);
+}
+
+// DEFAULT-LABEL: @test_scalbnf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_scalbnf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_scalbn_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_scalbnf(float x, int y) {
+  return scalbnf(x, y);
+}
+
+// DEFAULT-LABEL: @test_scalbn(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_scalbn(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_scalbn_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_scalbn(double x, int y) {
+  return scalbn(x, y);
+}
+
+// CHECK-LABEL: @test___signbitf(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// CHECK-NEXT:    [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
+// CHECK-NEXT:    [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
+// CHECK-NEXT:    ret i32 [[CONV]]
+//
+extern "C" __device__ BOOL_TYPE test___signbitf(float x) {
+  return __signbitf(x);
+}
+
+// CHECK-LABEL: @test___signbit(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// CHECK-NEXT:    [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
+// CHECK-NEXT:    [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
+// CHECK-NEXT:    ret i32 [[CONV]]
+//
+extern "C" __device__ BOOL_TYPE test___signbit(double x) {
+  return __signbit(x);
+}
+
+// DEFAULT-LABEL: @test_sincosf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
+// DEFAULT-NEXT:    call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = call contract float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// DEFAULT-NEXT:    store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]]
+// DEFAULT-NEXT:    [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]]
+// DEFAULT-NEXT:    store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]]
+// DEFAULT-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT:    ret void
+//
+// FINITEONLY-LABEL: @test_sincosf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
+// FINITEONLY-NEXT:    call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]]
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]]
+// FINITEONLY-NEXT:    store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]]
+// FINITEONLY-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT:    ret void
+//
+extern "C" __device__ void test_sincosf(float x, float *y, float *z) {
+  sincosf(x, y, z);
+}
+
+// DEFAULT-LABEL: @test_sincos(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
+// DEFAULT-NEXT:    call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = call contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// DEFAULT-NEXT:    store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]]
+// DEFAULT-NEXT:    [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15]]
+// DEFAULT-NEXT:    store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA15]]
+// DEFAULT-NEXT:    call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT:    ret void
+//
+// FINITEONLY-LABEL: @test_sincos(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
+// FINITEONLY-NEXT:    call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]]
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15]]
+// FINITEONLY-NEXT:    store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA15]]
+// FINITEONLY-NEXT:    call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT:    ret void
+//
+extern "C" __device__ void test_sincos(double x, double *y, double *z) {
+  sincos(x, y, z);
+}
+
+// DEFAULT-LABEL: @test_sincospif(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
+// DEFAULT-NEXT:    call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = call contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// DEFAULT-NEXT:    store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]]
+// DEFAULT-NEXT:    [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]]
+// DEFAULT-NEXT:    store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]]
+// DEFAULT-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT:    ret void
+//
+// FINITEONLY-LABEL: @test_sincospif(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
+// FINITEONLY-NEXT:    call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = call nnan ninf contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]]
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA13]]
+// FINITEONLY-NEXT:    store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]]
+// FINITEONLY-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT:    ret void
+//
+extern "C" __device__ void test_sincospif(float x, float *y, float *z) {
+  sincospif(x, y, z);
+}
+
+// DEFAULT-LABEL: @test_sincospi(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
+// DEFAULT-NEXT:    call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = call contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// DEFAULT-NEXT:    store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]]
+// DEFAULT-NEXT:    [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15]]
+// DEFAULT-NEXT:    store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA15]]
+// DEFAULT-NEXT:    call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// DEFAULT-NEXT:    ret void
+//
+// FINITEONLY-LABEL: @test_sincospi(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
+// FINITEONLY-NEXT:    call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = call nnan ninf contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA15]]
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA15]]
+// FINITEONLY-NEXT:    store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA15]]
+// FINITEONLY-NEXT:    call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// FINITEONLY-NEXT:    ret void
+//
+extern "C" __device__ void test_sincospi(double x, double *y, double *z) {
+  sincospi(x, y, z);
+}
+
+// DEFAULT-LABEL: @test_sinf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_sinf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_sinf(float x) {
+  return sinf(x);
+}
+
+// DEFAULT-LABEL: @test_sin(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sin_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_sin(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sin_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_sin(double x) {
+  return sin(x);
+}
+
+// DEFAULT-LABEL: @test_sinpif(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sinpi_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_sinpif(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sinpi_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_sinpif(float x) {
+  return sinpif(x);
+}
+
+// DEFAULT-LABEL: @test_sinpi(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sinpi_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_sinpi(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sinpi_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_sinpi(double x) {
+  return sinpi(x);
+}
+
+// DEFAULT-LABEL: @test_sqrtf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_sqrtf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_sqrtf(float x) {
+  return sqrtf(x);
+}
+
+// DEFAULT-LABEL: @test_sqrt(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_sqrt(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_sqrt(double x) {
+  return sqrt(x);
+}
+
+// DEFAULT-LABEL: @test_tanf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_tanf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_tanf(float x) {
+  return tanf(x);
+}
+
+// DEFAULT-LABEL: @test_tan(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_tan_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_tan(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_tan_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_tan(double x) {
+  return tan(x);
+}
+
+// DEFAULT-LABEL: @test_tanhf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_tanh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_tanhf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tanh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_tanhf(float x) {
+  return tanhf(x);
+}
+
+// DEFAULT-LABEL: @test_tanh(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_tanh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_tanh(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_tanh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_tanh(double x) {
+  return tanh(x);
+}
+
+// DEFAULT-LABEL: @test_tgammaf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_tgamma_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_tgammaf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tgamma_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_tgammaf(float x) {
+  return tgammaf(x);
+}
+
+// DEFAULT-LABEL: @test_tgamma(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_tgamma_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_tgamma(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_tgamma_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_tgamma(double x) {
+  return tgamma(x);
+}
+
+// DEFAULT-LABEL: @test_truncf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_trunc_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_truncf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_trunc_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_truncf(float x) {
+  return truncf(x);
+}
+
+// DEFAULT-LABEL: @test_trunc(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_trunc_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_trunc(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_trunc_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_trunc(double x) {
+  return trunc(x);
+}
+
+// DEFAULT-LABEL: @test_y0f(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_y0f(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_y0f(float x) {
+  return y0f(x);
+}
+
+// DEFAULT-LABEL: @test_y0(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_y0(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_y0(double x) {
+  return y0(x);
+}
+
+// DEFAULT-LABEL: @test_y1f(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_y1f(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test_y1f(float x) {
+  return y1f(x);
+}
+
+// DEFAULT-LABEL: @test_y1(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test_y1(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test_y1(double x) {
+  return y1(x);
+}
+
+// DEFAULT-LABEL: @test_ynf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
+// DEFAULT-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
+// DEFAULT-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
+// DEFAULT-NEXT:    ]
+// DEFAULT:       if.then.i:
+// DEFAULT-NEXT:    [[CALL_I_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    br label [[_ZL3YNFIF_EXIT:%.*]]
+// DEFAULT:       if.then2.i:
+// DEFAULT-NEXT:    [[CALL_I20_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT:    br label [[_ZL3YNFIF_EXIT]]
+// DEFAULT:       if.end4.i:
+// DEFAULT-NEXT:    [[CALL_I21_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT:    [[CALL_I22_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
+// DEFAULT:       for.body.i:
+// DEFAULT-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
+// DEFAULT-NEXT:    [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
+// DEFAULT-NEXT:    [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]]
+// DEFAULT-NEXT:    [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]]
+// DEFAULT-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// DEFAULT-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
+// DEFAULT:       _ZL3ynfif.exit:
+// 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]] ]
+// DEFAULT-NEXT:    ret float [[RETVAL_0_I]]
+//
+// FINITEONLY-LABEL: @test_ynf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
+// FINITEONLY-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
+// FINITEONLY-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
+// FINITEONLY-NEXT:    ]
+// FINITEONLY:       if.then.i:
+// FINITEONLY-NEXT:    [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    br label [[_ZL3YNFIF_EXIT:%.*]]
+// FINITEONLY:       if.then2.i:
+// FINITEONLY-NEXT:    [[CALL_I20_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    br label [[_ZL3YNFIF_EXIT]]
+// FINITEONLY:       if.end4.i:
+// FINITEONLY-NEXT:    [[CALL_I21_I:%.*]] = tail call nnan ninf contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    [[CALL_I22_I:%.*]] = tail call nnan ninf contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
+// FINITEONLY:       for.body.i:
+// FINITEONLY-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
+// FINITEONLY-NEXT:    [[DIV_I:%.*]] = fdiv nnan ninf contract float [[CONV_I]], [[Y]]
+// FINITEONLY-NEXT:    [[MUL8_I:%.*]] = fmul nnan ninf contract float [[__X1_025_I]], [[DIV_I]]
+// FINITEONLY-NEXT:    [[SUB_I]] = fsub nnan ninf contract float [[MUL8_I]], [[__X0_024_I]]
+// FINITEONLY-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// FINITEONLY-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
+// FINITEONLY:       _ZL3ynfif.exit:
+// 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]] ]
+// FINITEONLY-NEXT:    ret float [[RETVAL_0_I]]
+//
+extern "C" __device__ float test_ynf(int x, float y) {
+  return ynf(x, y);
+}
+
+// DEFAULT-LABEL: @test_yn(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
+// DEFAULT-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
+// DEFAULT-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
+// DEFAULT-NEXT:    ]
+// DEFAULT:       if.then.i:
+// DEFAULT-NEXT:    [[CALL_I_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    br label [[_ZL2YNID_EXIT:%.*]]
+// DEFAULT:       if.then2.i:
+// DEFAULT-NEXT:    [[CALL_I20_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT:    br label [[_ZL2YNID_EXIT]]
+// DEFAULT:       if.end4.i:
+// DEFAULT-NEXT:    [[CALL_I21_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT:    [[CALL_I22_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]]
+// DEFAULT-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// DEFAULT-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
+// DEFAULT:       for.body.i:
+// DEFAULT-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// DEFAULT-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
+// DEFAULT-NEXT:    [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
+// DEFAULT-NEXT:    [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]]
+// DEFAULT-NEXT:    [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]]
+// DEFAULT-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
+// DEFAULT-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// DEFAULT-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
+// DEFAULT:       _ZL2ynid.exit:
+// 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]] ]
+// DEFAULT-NEXT:    ret double [[RETVAL_0_I]]
+//
+// FINITEONLY-LABEL: @test_yn(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
+// FINITEONLY-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
+// FINITEONLY-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
+// FINITEONLY-NEXT:    ]
+// FINITEONLY:       if.then.i:
+// FINITEONLY-NEXT:    [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    br label [[_ZL2YNID_EXIT:%.*]]
+// FINITEONLY:       if.then2.i:
+// FINITEONLY-NEXT:    [[CALL_I20_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    br label [[_ZL2YNID_EXIT]]
+// FINITEONLY:       if.end4.i:
+// FINITEONLY-NEXT:    [[CALL_I21_I:%.*]] = tail call nnan ninf contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    [[CALL_I22_I:%.*]] = tail call nnan ninf contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// FINITEONLY-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
+// FINITEONLY:       for.body.i:
+// FINITEONLY-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// FINITEONLY-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
+// FINITEONLY-NEXT:    [[DIV_I:%.*]] = fdiv nnan ninf contract double [[CONV_I]], [[Y]]
+// FINITEONLY-NEXT:    [[MUL8_I:%.*]] = fmul nnan ninf contract double [[__X1_025_I]], [[DIV_I]]
+// FINITEONLY-NEXT:    [[SUB_I]] = fsub nnan ninf contract double [[MUL8_I]], [[__X0_024_I]]
+// FINITEONLY-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
+// FINITEONLY-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// FINITEONLY-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
+// FINITEONLY:       _ZL2ynid.exit:
+// 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]] ]
+// FINITEONLY-NEXT:    ret double [[RETVAL_0_I]]
+//
+extern "C" __device__ double test_yn(int x, double y) {
+  return yn(x, y);
+}
+
+// DEFAULT-LABEL: @test___cosf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test___cosf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___cosf(float x) {
+  return __cosf(x);
+}
+
+// DEFAULT-LABEL: @test___exp10f(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_native_exp10_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test___exp10f(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_exp10_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___exp10f(float x) {
+  return __exp10f(x);
+}
+
+// DEFAULT-LABEL: @test___expf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_native_exp_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test___expf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_exp_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___expf(float x) {
+  return __expf(x);
+}
+
+// DEFAULT-LABEL: @test___fadd_rn(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[ADD_I:%.*]] = fadd contract float [[X:%.*]], [[Y:%.*]]
+// DEFAULT-NEXT:    ret float [[ADD_I]]
+//
+// FINITEONLY-LABEL: @test___fadd_rn(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[ADD_I:%.*]] = fadd nnan ninf contract float [[X:%.*]], [[Y:%.*]]
+// FINITEONLY-NEXT:    ret float [[ADD_I]]
+//
+extern "C" __device__ float test___fadd_rn(float x, float y) {
+  return __fadd_rn(x, y);
+}
+
+// DEFAULT-LABEL: @test___fdividef(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[DIV_I:%.*]] = fdiv contract float [[X:%.*]], [[Y:%.*]]
+// DEFAULT-NEXT:    ret float [[DIV_I]]
+//
+// FINITEONLY-LABEL: @test___fdividef(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[DIV_I:%.*]] = fdiv nnan ninf contract float [[X:%.*]], [[Y:%.*]]
+// FINITEONLY-NEXT:    ret float [[DIV_I]]
+//
+extern "C" __device__ float test___fdividef(float x, float y) {
+  return __fdividef(x, y);
+}
+
+// DEFAULT-LABEL: @test__fmaf_rn(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test__fmaf_rn(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_fma_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test__fmaf_rn(float x, float y, float z) {
+  return __fmaf_rn(x, y, z);
+}
+
+// DEFAULT-LABEL: @test___fmul_rn(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[MUL_I:%.*]] = fmul contract float [[X:%.*]], [[Y:%.*]]
+// DEFAULT-NEXT:    ret float [[MUL_I]]
+//
+// FINITEONLY-LABEL: @test___fmul_rn(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[MUL_I:%.*]] = fmul nnan ninf contract float [[X:%.*]], [[Y:%.*]]
+// FINITEONLY-NEXT:    ret float [[MUL_I]]
+//
+extern "C" __device__ float test___fmul_rn(float x, float y) {
+  return __fmul_rn(x, y);
+}
+
+// DEFAULT-LABEL: @test___frcp_rn(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[DIV_I:%.*]] = fdiv contract float 1.000000e+00, [[X:%.*]]
+// DEFAULT-NEXT:    ret float [[DIV_I]]
+//
+// FINITEONLY-LABEL: @test___frcp_rn(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[DIV_I:%.*]] = fdiv nnan ninf contract float 1.000000e+00, [[X:%.*]]
+// FINITEONLY-NEXT:    ret float [[DIV_I]]
+//
+extern "C" __device__ float test___frcp_rn(float x) {
+  return __frcp_rn(x);
+}
+
+// DEFAULT-LABEL: @test___fsqrt_rn(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_native_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test___fsqrt_rn(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___fsqrt_rn(float x) {
+  return __fsqrt_rn(x);
+}
+
+// DEFAULT-LABEL: @test___fsub_rn(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[SUB_I:%.*]] = fsub contract float [[X:%.*]], [[Y:%.*]]
+// DEFAULT-NEXT:    ret float [[SUB_I]]
+//
+// FINITEONLY-LABEL: @test___fsub_rn(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[SUB_I:%.*]] = fsub nnan ninf contract float [[X:%.*]], [[Y:%.*]]
+// FINITEONLY-NEXT:    ret float [[SUB_I]]
+//
+extern "C" __device__ float test___fsub_rn(float x, float y) {
+  return __fsub_rn(x, y);
+}
+
+// DEFAULT-LABEL: @test___log10f(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_native_log10_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test___log10f(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_log10_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___log10f(float x) {
+  return __log10f(x);
+}
+
+// DEFAULT-LABEL: @test___log2f(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_native_log2_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test___log2f(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_log2_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___log2f(float x) {
+  return __log2f(x);
+}
+
+// DEFAULT-LABEL: @test___logf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_native_log_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test___logf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_log_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___logf(float x) {
+  return __logf(x);
+}
+
+// DEFAULT-LABEL: @test___powf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test___powf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___powf(float x, float y) {
+  return __powf(x, y);
+}
+
+// DEFAULT-LABEL: @test___saturatef(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CMP_I:%.*]] = fcmp contract olt float [[X:%.*]], 0.000000e+00
+// DEFAULT-NEXT:    [[CMP1_I:%.*]] = fcmp contract ogt float [[X]], 1.000000e+00
+// DEFAULT-NEXT:    [[COND_I:%.*]] = select contract i1 [[CMP1_I]], float 1.000000e+00, float [[X]]
+// DEFAULT-NEXT:    [[COND5_I:%.*]] = select contract i1 [[CMP_I]], float 0.000000e+00, float [[COND_I]]
+// DEFAULT-NEXT:    ret float [[COND5_I]]
+//
+// FINITEONLY-LABEL: @test___saturatef(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CMP_I:%.*]] = fcmp nnan ninf contract olt float [[X:%.*]], 0.000000e+00
+// FINITEONLY-NEXT:    [[CMP1_I:%.*]] = fcmp nnan ninf contract ogt float [[X]], 1.000000e+00
+// FINITEONLY-NEXT:    [[COND_I:%.*]] = select nnan ninf contract i1 [[CMP1_I]], float 1.000000e+00, float [[X]]
+// FINITEONLY-NEXT:    [[COND5_I:%.*]] = select nnan ninf contract i1 [[CMP_I]], float 0.000000e+00, float [[COND_I]]
+// FINITEONLY-NEXT:    ret float [[COND5_I]]
+//
+extern "C" __device__ float test___saturatef(float x) {
+  return __saturatef(x);
+}
+
+// DEFAULT-LABEL: @test___sincosf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]]
+// DEFAULT-NEXT:    [[CALL1_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]]
+// DEFAULT-NEXT:    store float [[CALL1_I]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]]
+// DEFAULT-NEXT:    ret void
+//
+// FINITEONLY-LABEL: @test___sincosf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA13]]
+// FINITEONLY-NEXT:    [[CALL1_I:%.*]] = tail call nnan ninf contract float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    store float [[CALL1_I]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA13]]
+// FINITEONLY-NEXT:    ret void
+//
+extern "C" __device__ void test___sincosf(float x, float *y, float *z) {
+  __sincosf(x, y, z);
+}
+
+// DEFAULT-LABEL: @test___sinf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test___sinf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___sinf(float x) {
+  return __sinf(x);
+}
+
+// DEFAULT-LABEL: @test___tanf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    ret float [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test___tanf(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___tanf(float x) {
+  return __tanf(x);
+}
+
+// DEFAULT-LABEL: @test___dadd_rn(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[ADD_I:%.*]] = fadd contract double [[X:%.*]], [[Y:%.*]]
+// DEFAULT-NEXT:    ret double [[ADD_I]]
+//
+// FINITEONLY-LABEL: @test___dadd_rn(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[ADD_I:%.*]] = fadd nnan ninf contract double [[X:%.*]], [[Y:%.*]]
+// FINITEONLY-NEXT:    ret double [[ADD_I]]
+//
+extern "C" __device__ double test___dadd_rn(double x, double y) {
+  return __dadd_rn(x, y);
+}
+
+// DEFAULT-LABEL: @test___ddiv_rn(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[DIV_I:%.*]] = fdiv contract double [[X:%.*]], [[Y:%.*]]
+// DEFAULT-NEXT:    ret double [[DIV_I]]
+//
+// FINITEONLY-LABEL: @test___ddiv_rn(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[DIV_I:%.*]] = fdiv nnan ninf contract double [[X:%.*]], [[Y:%.*]]
+// FINITEONLY-NEXT:    ret double [[DIV_I]]
+//
+extern "C" __device__ double test___ddiv_rn(double x, double y) {
+  return __ddiv_rn(x, y);
+}
+
+// DEFAULT-LABEL: @test___dmul_rn(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[MUL_I:%.*]] = fmul contract double [[X:%.*]], [[Y:%.*]]
+// DEFAULT-NEXT:    ret double [[MUL_I]]
+//
+// FINITEONLY-LABEL: @test___dmul_rn(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[MUL_I:%.*]] = fmul nnan ninf contract double [[X:%.*]], [[Y:%.*]]
+// FINITEONLY-NEXT:    ret double [[MUL_I]]
+//
+extern "C" __device__ double test___dmul_rn(double x, double y) {
+  return __dmul_rn(x, y);
+}
+
+// DEFAULT-LABEL: @test___drcp_rn(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[DIV_I:%.*]] = fdiv contract double 1.000000e+00, [[X:%.*]]
+// DEFAULT-NEXT:    ret double [[DIV_I]]
+//
+// FINITEONLY-LABEL: @test___drcp_rn(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[DIV_I:%.*]] = fdiv nnan ninf contract double 1.000000e+00, [[X:%.*]]
+// FINITEONLY-NEXT:    ret double [[DIV_I]]
+//
+extern "C" __device__ double test___drcp_rn(double x) {
+  return __drcp_rn(x);
+}
+
+// DEFAULT-LABEL: @test___dsqrt_rn(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test___dsqrt_rn(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test___dsqrt_rn(double x) {
+  return __dsqrt_rn(x);
+}
+
+// DEFAULT-LABEL: @test__fma_rn(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I]]
+//
+// FINITEONLY-LABEL: @test__fma_rn(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract double @__ocml_fma_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test__fma_rn(double x, double y, double z) {
+  return __fma_rn(x, y, z);
+}
+
+// DEFAULT-LABEL: @test_float_min(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I_I]]
+//
+// FINITEONLY-LABEL: @test_float_min(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_fmin_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I_I]]
+//
+extern "C" __device__ float test_float_min(float x, float y) {
+  return min(x, y);
+}
+
+// DEFAULT-LABEL: @test_float_max(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I_I:%.*]] = tail call contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret float [[CALL_I_I]]
+//
+// FINITEONLY-LABEL: @test_float_max(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I_I:%.*]] = tail call nnan ninf contract float @__ocml_fmax_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret float [[CALL_I_I]]
+//
+extern "C" __device__ float test_float_max(float x, float y) {
+  return max(x, y);
+}
+
+// DEFAULT-LABEL: @test_double_min(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I_I:%.*]] = tail call contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I_I]]
+//
+// FINITEONLY-LABEL: @test_double_min(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_fmin_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I_I]]
+//
+extern "C" __device__ double test_double_min(double x, double y) {
+  return min(x, y);
+}
+
+// DEFAULT-LABEL: @test_double_max(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I_I:%.*]] = tail call contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// DEFAULT-NEXT:    ret double [[CALL_I_I]]
+//
+// FINITEONLY-LABEL: @test_double_max(
+// FINITEONLY-NEXT:  entry:
+// FINITEONLY-NEXT:    [[CALL_I_I:%.*]] = tail call nnan ninf contract double @__ocml_fmax_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// FINITEONLY-NEXT:    ret double [[CALL_I_I]]
+//
+extern "C" __device__ double test_double_max(double x, double y) {
+  return max(x, y);
+}
+// CHECK-LABEL: @test_int_min(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[TMP0:%.*]] = tail call i32 @llvm.smin.i32(i32 [[X:%.*]], i32 [[Y:%.*]])
+// CHECK-NEXT:    ret i32 [[TMP0]]
+//
+extern "C" __device__ int test_int_min(int x, int y) {
+  return min(x, y);
+}
+
+// CHECK-LABEL: @test_int_max(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[TMP0:%.*]] = tail call i32 @llvm.smax.i32(i32 [[X:%.*]], i32 [[Y:%.*]])
+// CHECK-NEXT:    ret i32 [[TMP0]]
+//
+extern "C" __device__ int test_int_max(int x, int y) {
+  return max(x, y);
+}

clang/test/Headers/__clang_hip_math_ocml_rounded_ops.hip

+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
+// Test with OCML_BASIC_ROUNDED_OPERATIONS
+// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h      \
+// RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
+// RUN:   -internal-isystem %S/Inputs/include \
+// RUN:   -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
+// RUN:   -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -O1 -o - \
+// RUN:   -D__HIPCC_RTC__ -DOCML_BASIC_ROUNDED_OPERATIONS | FileCheck %s
+
+// CHECK-LABEL: @test___fadd_rd(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_add_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4:[0-9]+]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___fadd_rd(float x, float y) {
+  return __fadd_rd(x, y);
+}
+
+// CHECK-LABEL: @test___fadd_rn(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_add_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___fadd_rn(float x, float y) {
+  return __fadd_rn(x, y);
+}
+
+// CHECK-LABEL: @test___fadd_ru(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_add_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___fadd_ru(float x, float y) {
+  return __fadd_ru(x, y);
+}
+
+// CHECK-LABEL: @test___fadd_rz(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_add_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___fadd_rz(float x, float y) {
+  return __fadd_rz(x, y);
+}
+
+// CHECK-LABEL: @test__fmaf_rd(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test__fmaf_rd(float x, float y, float z) {
+  return __fmaf_rd(x, y, z);
+}
+
+// CHECK-LABEL: @test__fmaf_rn(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test__fmaf_rn(float x, float y, float z) {
+  return __fmaf_rn(x, y, z);
+}
+
+// CHECK-LABEL: @test__fmaf_ru(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test__fmaf_ru(float x, float y, float z) {
+  return __fmaf_ru(x, y, z);
+}
+
+// CHECK-LABEL: @test__fmaf_rz(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_fma_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test__fmaf_rz(float x, float y, float z) {
+  return __fmaf_rz(x, y, z);
+}
+
+// CHECK-LABEL: @test___fmul_rd(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___fmul_rd(float x, float y) {
+  return __fmul_rd(x, y);
+}
+
+// CHECK-LABEL: @test___fmul_rn(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___fmul_rn(float x, float y) {
+  return __fmul_rn(x, y);
+}
+
+// CHECK-LABEL: @test___fmul_ru(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___fmul_ru(float x, float y) {
+  return __fmul_ru(x, y);
+}
+
+// CHECK-LABEL: @test___fmul_rz(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_mul_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___fmul_rz(float x, float y) {
+  return __fmul_rz(x, y);
+}
+
+// CHECK-LABEL: @test___frcp_rd(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_div_rtn_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___frcp_rd(float x) {
+  return __frcp_rd(x);
+}
+
+// CHECK-LABEL: @test___frcp_rn(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_div_rte_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___frcp_rn(float x) {
+  return __frcp_rn(x);
+}
+
+// CHECK-LABEL: @test___frcp_ru(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_div_rtp_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___frcp_ru(float x) {
+  return __frcp_ru(x);
+}
+
+// CHECK-LABEL: @test___frcp_rz(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_div_rtz_f32(float noundef 1.000000e+00, float noundef [[X:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___frcp_rz(float x) {
+  return __frcp_rz(x);
+}
+
+// CHECK-LABEL: @test___fsqrt_rd(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rtn_f32(float noundef [[X:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___fsqrt_rd(float x) {
+  return __fsqrt_rd(x);
+}
+
+// CHECK-LABEL: @test___fsqrt_rn(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rte_f32(float noundef [[X:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___fsqrt_rn(float x) {
+  return __fsqrt_rn(x);
+}
+
+// CHECK-LABEL: @test___fsqrt_ru(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rtp_f32(float noundef [[X:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___fsqrt_ru(float x) {
+  return __fsqrt_ru(x);
+}
+
+// CHECK-LABEL: @test___fsqrt_rz(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_rtz_f32(float noundef [[X:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___fsqrt_rz(float x) {
+  return __fsqrt_rz(x);
+}
+
+// CHECK-LABEL: @test___fsub_rd(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rtn_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___fsub_rd(float x, float y) {
+  return __fsub_rd(x, y);
+}
+
+// CHECK-LABEL: @test___fsub_rn(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rte_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___fsub_rn(float x, float y) {
+  return __fsub_rn(x, y);
+}
+
+// CHECK-LABEL: @test___fsub_ru(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rtp_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___fsub_ru(float x, float y) {
+  return __fsub_ru(x, y);
+}
+
+// CHECK-LABEL: @test___fsub_rz(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sub_rtz_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret float [[CALL_I]]
+//
+extern "C" __device__ float test___fsub_rz(float x, float y) {
+  return __fsub_rz(x, y);
+}
+
+// CHECK-LABEL: @test___dadd_rd(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_add_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test___dadd_rd(double x, double y) {
+  return __dadd_rd(x, y);
+}
+
+// CHECK-LABEL: @test___dadd_rn(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_add_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test___dadd_rn(double x, double y) {
+  return __dadd_rn(x, y);
+}
+
+// CHECK-LABEL: @test___dadd_ru(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_add_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test___dadd_ru(double x, double y) {
+  return __dadd_ru(x, y);
+}
+
+// CHECK-LABEL: @test___dadd_rz(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_add_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test___dadd_rz(double x, double y) {
+  return __dadd_rz(x, y);
+}
+
+// CHECK-LABEL: @test___dmul_rd(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test___dmul_rd(double x, double y) {
+  return __dmul_rd(x, y);
+}
+
+// CHECK-LABEL: @test___dmul_rn(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test___dmul_rn(double x, double y) {
+  return __dmul_rn(x, y);
+}
+
+// CHECK-LABEL: @test___dmul_ru(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test___dmul_ru(double x, double y) {
+  return __dmul_ru(x, y);
+}
+
+// CHECK-LABEL: @test___dmul_rz(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_mul_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test___dmul_rz(double x, double y) {
+  return __dmul_rz(x, y);
+}
+
+// CHECK-LABEL: @test___drcp_rd(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CONV:%.*]] = fpext float [[X:%.*]] to double
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_div_rtn_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR4]]
+// CHECK-NEXT:    [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
+// CHECK-NEXT:    ret float [[CONV1]]
+//
+extern "C" __device__ float test___drcp_rd(float x) {
+  return __drcp_rd(x);
+}
+
+// CHECK-LABEL: @test___drcp_rn(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CONV:%.*]] = fpext float [[X:%.*]] to double
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_div_rte_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR4]]
+// CHECK-NEXT:    [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
+// CHECK-NEXT:    ret float [[CONV1]]
+//
+extern "C" __device__ float test___drcp_rn(float x) {
+  return __drcp_rn(x);
+}
+
+// CHECK-LABEL: @test___drcp_ru(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CONV:%.*]] = fpext float [[X:%.*]] to double
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_div_rtp_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR4]]
+// CHECK-NEXT:    [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
+// CHECK-NEXT:    ret float [[CONV1]]
+//
+extern "C" __device__ float test___drcp_ru(float x) {
+  return __drcp_ru(x);
+}
+
+// CHECK-LABEL: @test___drcp_rz(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CONV:%.*]] = fpext float [[X:%.*]] to double
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_div_rtz_f64(double noundef 1.000000e+00, double noundef [[CONV]]) #[[ATTR4]]
+// CHECK-NEXT:    [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
+// CHECK-NEXT:    ret float [[CONV1]]
+//
+extern "C" __device__ float test___drcp_rz(float x) {
+  return __drcp_rz(x);
+}
+
+// CHECK-LABEL: @test___dsqrt_rd(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CONV:%.*]] = fpext float [[X:%.*]] to double
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rtn_f64(double noundef [[CONV]]) #[[ATTR4]]
+// CHECK-NEXT:    [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
+// CHECK-NEXT:    ret float [[CONV1]]
+//
+extern "C" __device__ float test___dsqrt_rd(float x) {
+  return __dsqrt_rd(x);
+}
+
+// CHECK-LABEL: @test___dsqrt_rn(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CONV:%.*]] = fpext float [[X:%.*]] to double
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rte_f64(double noundef [[CONV]]) #[[ATTR4]]
+// CHECK-NEXT:    [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
+// CHECK-NEXT:    ret float [[CONV1]]
+//
+extern "C" __device__ float test___dsqrt_rn(float x) {
+  return __dsqrt_rn(x);
+}
+
+// CHECK-LABEL: @test___dsqrt_ru(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CONV:%.*]] = fpext float [[X:%.*]] to double
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rtp_f64(double noundef [[CONV]]) #[[ATTR4]]
+// CHECK-NEXT:    [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
+// CHECK-NEXT:    ret float [[CONV1]]
+//
+extern "C" __device__ float test___dsqrt_ru(float x) {
+  return __dsqrt_ru(x);
+}
+
+// CHECK-LABEL: @test___dsqrt_rz(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CONV:%.*]] = fpext float [[X:%.*]] to double
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_rtz_f64(double noundef [[CONV]]) #[[ATTR4]]
+// CHECK-NEXT:    [[CONV1:%.*]] = fptrunc double [[CALL_I]] to float
+// CHECK-NEXT:    ret float [[CONV1]]
+//
+extern "C" __device__ float test___dsqrt_rz(float x) {
+  return __dsqrt_rz(x);
+}
+
+// CHECK-LABEL: @test__fma_rd(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rtn_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test__fma_rd(double x, double y, double z) {
+  return __fma_rd(x, y, z);
+}
+
+// CHECK-LABEL: @test__fma_rn(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rte_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test__fma_rn(double x, double y, double z) {
+  return __fma_rn(x, y, z);
+}
+
+// CHECK-LABEL: @test__fma_ru(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rtp_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test__fma_ru(double x, double y, double z) {
+  return __fma_ru(x, y, z);
+}
+
+// CHECK-LABEL: @test__fma_rz(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_fma_rtz_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR4]]
+// CHECK-NEXT:    ret double [[CALL_I]]
+//
+extern "C" __device__ double test__fma_rz(double x, double y, double z) {
+  return __fma_rz(x, y, z);
+}