[clang] a0709e2 - [HIP] Use native math functions for `-fcuda-approx-transcendentals`

[Yaxun Liu via cfe-commits]

Author: Yaxun (Sam) Liu
Date: 2023-07-25T10:04:31-04:00
New Revision: a0709e226b482c3d4a2d1f15227641590c43527a

URL: https://github.com/llvm/llvm-project/commit/a0709e226b482c3d4a2d1f15227641590c43527a
DIFF: https://github.com/llvm/llvm-project/commit/a0709e226b482c3d4a2d1f15227641590c43527a.diff

LOG: [HIP] Use native math functions for `-fcuda-approx-transcendentals`

CUDA allows replacing standard math functions with less accurate
native math functions when -use_fast_math is specified
(https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#intrinsic-functions).
Cuda-clang does this when -fcuda-approx-transcendentals
or -ffast-math is specified. HIP-clang currently passes option
-fcuda-approx-transcendentals to clang -cc1 and predefines
__CLANG_CUDA_APPROX_TRANSCENDENTALS__ but does
not replace standard math functions with native math functions.

This patch implements this in a similar approach as cuda-clang.

Reviewed by: Brian Sumner, Matt Arsenault

Differential Revision: https://reviews.llvm.org/D154790

Added: 
    

Modified: 
    clang/lib/Headers/__clang_hip_math.h
    clang/test/Headers/__clang_hip_math.hip

Removed: 
    


################################################################################
diff  --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h
index a47dda3327f438..7689f4221ce440 100644
--- a/clang/lib/Headers/__clang_hip_math.h
+++ b/clang/lib/Headers/__clang_hip_math.h
@@ -32,6 +32,17 @@
 #define __DEVICE__ static __device__ inline __attribute__((always_inline))
 #endif
 
+// Device library provides fast low precision and slow full-recision
+// implementations for some functions. Which one gets selected depends on
+// __CLANG_CUDA_APPROX_TRANSCENDENTALS__ which gets defined by clang if
+// -ffast-math or -fcuda-approx-transcendentals are in effect.
+#pragma push_macro("__FAST_OR_SLOW")
+#if defined(__CLANG_CUDA_APPROX_TRANSCENDENTALS__)
+#define __FAST_OR_SLOW(fast, slow) fast
+#else
+#define __FAST_OR_SLOW(fast, slow) slow
+#endif
+
 // A few functions return bool type starting only in C++11.
 #pragma push_macro("__RETURN_TYPE")
 #ifdef __OPENMP_AMDGCN__
@@ -139,6 +150,168 @@ uint64_t __make_mantissa(const char *__tagp __attribute__((nonnull))) {
 }
 
 // BEGIN FLOAT
+
+// BEGIN INTRINSICS
+
+__DEVICE__
+float __cosf(float __x) { return __ocml_native_cos_f32(__x); }
+
+__DEVICE__
+float __exp10f(float __x) {
+  const float __log2_10 = 0x1.a934f0p+1f;
+  return __builtin_amdgcn_exp2f(__log2_10 * __x);
+}
+
+__DEVICE__
+float __expf(float __x) {
+  const float __log2_e = 0x1.715476p+0;
+  return __builtin_amdgcn_exp2f(__log2_e * __x);
+}
+
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
+__DEVICE__
+float __fadd_rd(float __x, float __y) { return __ocml_add_rtn_f32(__x, __y); }
+__DEVICE__
+float __fadd_rn(float __x, float __y) { return __ocml_add_rte_f32(__x, __y); }
+__DEVICE__
+float __fadd_ru(float __x, float __y) { return __ocml_add_rtp_f32(__x, __y); }
+__DEVICE__
+float __fadd_rz(float __x, float __y) { return __ocml_add_rtz_f32(__x, __y); }
+#else
+__DEVICE__
+float __fadd_rn(float __x, float __y) { return __x + __y; }
+#endif
+
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
+__DEVICE__
+float __fdiv_rd(float __x, float __y) { return __ocml_div_rtn_f32(__x, __y); }
+__DEVICE__
+float __fdiv_rn(float __x, float __y) { return __ocml_div_rte_f32(__x, __y); }
+__DEVICE__
+float __fdiv_ru(float __x, float __y) { return __ocml_div_rtp_f32(__x, __y); }
+__DEVICE__
+float __fdiv_rz(float __x, float __y) { return __ocml_div_rtz_f32(__x, __y); }
+#else
+__DEVICE__
+float __fdiv_rn(float __x, float __y) { return __x / __y; }
+#endif
+
+__DEVICE__
+float __fdividef(float __x, float __y) { return __x / __y; }
+
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
+__DEVICE__
+float __fmaf_rd(float __x, float __y, float __z) {
+  return __ocml_fma_rtn_f32(__x, __y, __z);
+}
+__DEVICE__
+float __fmaf_rn(float __x, float __y, float __z) {
+  return __ocml_fma_rte_f32(__x, __y, __z);
+}
+__DEVICE__
+float __fmaf_ru(float __x, float __y, float __z) {
+  return __ocml_fma_rtp_f32(__x, __y, __z);
+}
+__DEVICE__
+float __fmaf_rz(float __x, float __y, float __z) {
+  return __ocml_fma_rtz_f32(__x, __y, __z);
+}
+#else
+__DEVICE__
+float __fmaf_rn(float __x, float __y, float __z) {
+  return __builtin_fmaf(__x, __y, __z);
+}
+#endif
+
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
+__DEVICE__
+float __fmul_rd(float __x, float __y) { return __ocml_mul_rtn_f32(__x, __y); }
+__DEVICE__
+float __fmul_rn(float __x, float __y) { return __ocml_mul_rte_f32(__x, __y); }
+__DEVICE__
+float __fmul_ru(float __x, float __y) { return __ocml_mul_rtp_f32(__x, __y); }
+__DEVICE__
+float __fmul_rz(float __x, float __y) { return __ocml_mul_rtz_f32(__x, __y); }
+#else
+__DEVICE__
+float __fmul_rn(float __x, float __y) { return __x * __y; }
+#endif
+
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
+__DEVICE__
+float __frcp_rd(float __x) { return __ocml_div_rtn_f32(1.0f, __x); }
+__DEVICE__
+float __frcp_rn(float __x) { return __ocml_div_rte_f32(1.0f, __x); }
+__DEVICE__
+float __frcp_ru(float __x) { return __ocml_div_rtp_f32(1.0f, __x); }
+__DEVICE__
+float __frcp_rz(float __x) { return __ocml_div_rtz_f32(1.0f, __x); }
+#else
+__DEVICE__
+float __frcp_rn(float __x) { return 1.0f / __x; }
+#endif
+
+__DEVICE__
+float __frsqrt_rn(float __x) { return __builtin_amdgcn_rsqf(__x); }
+
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
+__DEVICE__
+float __fsqrt_rd(float __x) { return __ocml_sqrt_rtn_f32(__x); }
+__DEVICE__
+float __fsqrt_rn(float __x) { return __ocml_sqrt_rte_f32(__x); }
+__DEVICE__
+float __fsqrt_ru(float __x) { return __ocml_sqrt_rtp_f32(__x); }
+__DEVICE__
+float __fsqrt_rz(float __x) { return __ocml_sqrt_rtz_f32(__x); }
+#else
+__DEVICE__
+float __fsqrt_rn(float __x) { return __ocml_native_sqrt_f32(__x); }
+#endif
+
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
+__DEVICE__
+float __fsub_rd(float __x, float __y) { return __ocml_sub_rtn_f32(__x, __y); }
+__DEVICE__
+float __fsub_rn(float __x, float __y) { return __ocml_sub_rte_f32(__x, __y); }
+__DEVICE__
+float __fsub_ru(float __x, float __y) { return __ocml_sub_rtp_f32(__x, __y); }
+__DEVICE__
+float __fsub_rz(float __x, float __y) { return __ocml_sub_rtz_f32(__x, __y); }
+#else
+__DEVICE__
+float __fsub_rn(float __x, float __y) { return __x - __y; }
+#endif
+
+__DEVICE__
+float __log10f(float __x) { return __builtin_log10f(__x); }
+
+__DEVICE__
+float __log2f(float __x) { return __builtin_amdgcn_logf(__x); }
+
+__DEVICE__
+float __logf(float __x) { return __builtin_logf(__x); }
+
+__DEVICE__
+float __powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); }
+
+__DEVICE__
+float __saturatef(float __x) { return (__x < 0) ? 0 : ((__x > 1) ? 1 : __x); }
+
+__DEVICE__
+void __sincosf(float __x, float *__sinptr, float *__cosptr) {
+  *__sinptr = __ocml_native_sin_f32(__x);
+  *__cosptr = __ocml_native_cos_f32(__x);
+}
+
+__DEVICE__
+float __sinf(float __x) { return __ocml_native_sin_f32(__x); }
+
+__DEVICE__
+float __tanf(float __x) {
+  return __sinf(__x) * __builtin_amdgcn_rcpf(__cosf(__x));
+}
+// END INTRINSICS
+
 #if defined(__cplusplus)
 __DEVICE__
 int abs(int __x) {
@@ -188,7 +361,7 @@ __DEVICE__
 float copysignf(float __x, float __y) { return __builtin_copysignf(__x, __y); }
 
 __DEVICE__
-float cosf(float __x) { return __ocml_cos_f32(__x); }
+float cosf(float __x) { return __FAST_OR_SLOW(__cosf, __ocml_cos_f32)(__x); }
 
 __DEVICE__
 float coshf(float __x) { return __ocml_cosh_f32(__x); }
@@ -321,13 +494,13 @@ __DEVICE__
 float log1pf(float __x) { return __ocml_log1p_f32(__x); }
 
 __DEVICE__
-float log2f(float __x) { return __builtin_log2f(__x); }
+float log2f(float __x) { return __FAST_OR_SLOW(__log2f, __ocml_log2_f32)(__x); }
 
 __DEVICE__
 float logbf(float __x) { return __ocml_logb_f32(__x); }
 
 __DEVICE__
-float logf(float __x) { return __builtin_logf(__x); }
+float logf(float __x) { return __FAST_OR_SLOW(__logf, __ocml_log_f32)(__x); }
 
 __DEVICE__
 long int lrintf(float __x) { return __builtin_rintf(__x); }
@@ -483,9 +656,13 @@ void sincosf(float __x, float *__sinptr, float *__cosptr) {
 #ifdef __OPENMP_AMDGCN__
 #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
 #endif
+#ifdef __CLANG_CUDA_APPROX_TRANSCENDENTALS__
+  __sincosf(__x, __sinptr, __cosptr);
+#else
   *__sinptr =
       __ocml_sincos_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
   *__cosptr = __tmp;
+#endif
 }
 
 __DEVICE__
@@ -500,7 +677,7 @@ void sincospif(float __x, float *__sinptr, float *__cosptr) {
 }
 
 __DEVICE__
-float sinf(float __x) { return __ocml_sin_f32(__x); }
+float sinf(float __x) { return __FAST_OR_SLOW(__sinf, __ocml_sin_f32)(__x); }
 
 __DEVICE__
 float sinhf(float __x) { return __ocml_sinh_f32(__x); }
@@ -551,158 +728,7 @@ float ynf(int __n, float __x) { // TODO: we could use Ahmes multiplication
   return __x1;
 }
 
-// BEGIN INTRINSICS
-
-__DEVICE__
-float __cosf(float __x) { return __ocml_native_cos_f32(__x); }
-
-__DEVICE__
-float __exp10f(float __x) { return __ocml_native_exp10_f32(__x); }
-
-__DEVICE__
-float __expf(float __x) { return __ocml_native_exp_f32(__x); }
-
-#if defined OCML_BASIC_ROUNDED_OPERATIONS
-__DEVICE__
-float __fadd_rd(float __x, float __y) { return __ocml_add_rtn_f32(__x, __y); }
-__DEVICE__
-float __fadd_rn(float __x, float __y) { return __ocml_add_rte_f32(__x, __y); }
-__DEVICE__
-float __fadd_ru(float __x, float __y) { return __ocml_add_rtp_f32(__x, __y); }
-__DEVICE__
-float __fadd_rz(float __x, float __y) { return __ocml_add_rtz_f32(__x, __y); }
-#else
-__DEVICE__
-float __fadd_rn(float __x, float __y) { return __x + __y; }
-#endif
-
-#if defined OCML_BASIC_ROUNDED_OPERATIONS
-__DEVICE__
-float __fdiv_rd(float __x, float __y) { return __ocml_div_rtn_f32(__x, __y); }
-__DEVICE__
-float __fdiv_rn(float __x, float __y) { return __ocml_div_rte_f32(__x, __y); }
-__DEVICE__
-float __fdiv_ru(float __x, float __y) { return __ocml_div_rtp_f32(__x, __y); }
-__DEVICE__
-float __fdiv_rz(float __x, float __y) { return __ocml_div_rtz_f32(__x, __y); }
-#else
-__DEVICE__
-float __fdiv_rn(float __x, float __y) { return __x / __y; }
-#endif
-
-__DEVICE__
-float __fdividef(float __x, float __y) { return __x / __y; }
-
-#if defined OCML_BASIC_ROUNDED_OPERATIONS
-__DEVICE__
-float __fmaf_rd(float __x, float __y, float __z) {
-  return __ocml_fma_rtn_f32(__x, __y, __z);
-}
-__DEVICE__
-float __fmaf_rn(float __x, float __y, float __z) {
-  return __ocml_fma_rte_f32(__x, __y, __z);
-}
-__DEVICE__
-float __fmaf_ru(float __x, float __y, float __z) {
-  return __ocml_fma_rtp_f32(__x, __y, __z);
-}
-__DEVICE__
-float __fmaf_rz(float __x, float __y, float __z) {
-  return __ocml_fma_rtz_f32(__x, __y, __z);
-}
-#else
-__DEVICE__
-float __fmaf_rn(float __x, float __y, float __z) {
-  return __builtin_fmaf(__x, __y, __z);
-}
-#endif
-
-#if defined OCML_BASIC_ROUNDED_OPERATIONS
-__DEVICE__
-float __fmul_rd(float __x, float __y) { return __ocml_mul_rtn_f32(__x, __y); }
-__DEVICE__
-float __fmul_rn(float __x, float __y) { return __ocml_mul_rte_f32(__x, __y); }
-__DEVICE__
-float __fmul_ru(float __x, float __y) { return __ocml_mul_rtp_f32(__x, __y); }
-__DEVICE__
-float __fmul_rz(float __x, float __y) { return __ocml_mul_rtz_f32(__x, __y); }
-#else
-__DEVICE__
-float __fmul_rn(float __x, float __y) { return __x * __y; }
-#endif
-
-#if defined OCML_BASIC_ROUNDED_OPERATIONS
-__DEVICE__
-float __frcp_rd(float __x) { return __ocml_div_rtn_f32(1.0f, __x); }
-__DEVICE__
-float __frcp_rn(float __x) { return __ocml_div_rte_f32(1.0f, __x); }
-__DEVICE__
-float __frcp_ru(float __x) { return __ocml_div_rtp_f32(1.0f, __x); }
-__DEVICE__
-float __frcp_rz(float __x) { return __ocml_div_rtz_f32(1.0f, __x); }
-#else
-__DEVICE__
-float __frcp_rn(float __x) { return 1.0f / __x; }
-#endif
-
-__DEVICE__
-float __frsqrt_rn(float __x) { return __builtin_amdgcn_rsqf(__x); }
-
-#if defined OCML_BASIC_ROUNDED_OPERATIONS
-__DEVICE__
-float __fsqrt_rd(float __x) { return __ocml_sqrt_rtn_f32(__x); }
-__DEVICE__
-float __fsqrt_rn(float __x) { return __ocml_sqrt_rte_f32(__x); }
-__DEVICE__
-float __fsqrt_ru(float __x) { return __ocml_sqrt_rtp_f32(__x); }
-__DEVICE__
-float __fsqrt_rz(float __x) { return __ocml_sqrt_rtz_f32(__x); }
-#else
-__DEVICE__
-float __fsqrt_rn(float __x) { return __ocml_native_sqrt_f32(__x); }
-#endif
-
-#if defined OCML_BASIC_ROUNDED_OPERATIONS
-__DEVICE__
-float __fsub_rd(float __x, float __y) { return __ocml_sub_rtn_f32(__x, __y); }
-__DEVICE__
-float __fsub_rn(float __x, float __y) { return __ocml_sub_rte_f32(__x, __y); }
-__DEVICE__
-float __fsub_ru(float __x, float __y) { return __ocml_sub_rtp_f32(__x, __y); }
-__DEVICE__
-float __fsub_rz(float __x, float __y) { return __ocml_sub_rtz_f32(__x, __y); }
-#else
-__DEVICE__
-float __fsub_rn(float __x, float __y) { return __x - __y; }
-#endif
-
-__DEVICE__
-float __log10f(float __x) { return __ocml_native_log10_f32(__x); }
-
-__DEVICE__
-float __log2f(float __x) { return __ocml_native_log2_f32(__x); }
-
-__DEVICE__
-float __logf(float __x) { return __ocml_native_log_f32(__x); }
-
-__DEVICE__
-float __powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); }
-
-__DEVICE__
-float __saturatef(float __x) { return (__x < 0) ? 0 : ((__x > 1) ? 1 : __x); }
-
-__DEVICE__
-void __sincosf(float __x, float *__sinptr, float *__cosptr) {
-  *__sinptr = __ocml_native_sin_f32(__x);
-  *__cosptr = __ocml_native_cos_f32(__x);
-}
-
-__DEVICE__
-float __sinf(float __x) { return __ocml_native_sin_f32(__x); }
 
-__DEVICE__
-float __tanf(float __x) { return __ocml_tan_f32(__x); }
-// END INTRINSICS
 // END FLOAT
 
 // BEGIN DOUBLE
@@ -1299,5 +1325,6 @@ __host__ inline static int max(int __arg1, int __arg2) {
 
 #pragma pop_macro("__DEVICE__")
 #pragma pop_macro("__RETURN_TYPE")
+#pragma pop_macro("__FAST_OR_SLOW")
 
 #endif // __CLANG_HIP_MATH_H__

diff  --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index 85d559f161da92..a4c00bfb4a1a04 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -17,6 +17,14 @@
 // 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
 
+// Check that we end up with -fapprox-func 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 -fcuda-approx-transcendentals -o - \
+// RUN:   -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,APPROX %s
+
 #define BOOL_TYPE int
 typedef unsigned long long uint64_t;
 
@@ -258,6 +266,11 @@ extern "C" __device__ long long test_llabs(long x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_acos_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14:[0-9]+]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_acosf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_acos_f32(float noundef [[X:%.*]]) #[[ATTR14:[0-9]+]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_acosf(float x) {
   return acosf(x);
 }
@@ -272,6 +285,11 @@ extern "C" __device__ float test_acosf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_acos_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_acos(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_acos_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_acos(double x) {
   return acos(x);
 }
@@ -286,6 +304,11 @@ extern "C" __device__ double test_acos(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_acosh_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15:[0-9]+]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_acoshf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_acosh_f32(float noundef [[X:%.*]]) #[[ATTR15:[0-9]+]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_acoshf(float x) {
   return acoshf(x);
 }
@@ -300,6 +323,11 @@ extern "C" __device__ float test_acoshf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_acosh_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_acosh(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_acosh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_acosh(double x) {
   return acosh(x);
 }
@@ -314,6 +342,11 @@ extern "C" __device__ double test_acosh(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_asin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_asinf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_asin_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_asinf(float x) {
   return asinf(x);
 }
@@ -328,6 +361,11 @@ extern "C" __device__ float test_asinf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_asin_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_asin(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_asin_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_asin(double x) {
 
   return asin(x);
@@ -343,6 +381,11 @@ extern "C" __device__ double test_asin(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_asinh_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_asinhf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_asinh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_asinhf(float x) {
   return asinhf(x);
 }
@@ -357,6 +400,11 @@ extern "C" __device__ float test_asinhf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_asinh_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_asinh(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_asinh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_asinh(double x) {
   return asinh(x);
 }
@@ -371,6 +419,11 @@ extern "C" __device__ double test_asinh(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_atan2_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_atan2f(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_atan2_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_atan2f(float x, float y) {
   return atan2f(x, y);
 }
@@ -385,6 +438,11 @@ extern "C" __device__ float test_atan2f(float x, float y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_atan2_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_atan2(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_atan2_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_atan2(double x, double y) {
   return atan2(x, y);
 }
@@ -399,6 +457,11 @@ extern "C" __device__ double test_atan2(double x, double y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_atan_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_atanf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_atan_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_atanf(float x) {
   return atanf(x);
 }
@@ -413,6 +476,11 @@ extern "C" __device__ float test_atanf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_atan_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_atan(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_atan_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_atan(double x) {
   return atan(x);
 }
@@ -427,6 +495,11 @@ extern "C" __device__ double test_atan(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_atanh_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_atanhf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_atanh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_atanhf(float x) {
   return atanhf(x);
 }
@@ -441,6 +514,11 @@ extern "C" __device__ float test_atanhf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_atanh_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_atanh(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_atanh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_atanh(double x) {
   return atanh(x);
 }
@@ -455,6 +533,11 @@ extern "C" __device__ double test_atanh(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_cbrt_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_cbrtf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_cbrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_cbrtf(float x) {
   return cbrtf(x);
 }
@@ -469,6 +552,11 @@ extern "C" __device__ float test_cbrtf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_cbrt_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_cbrt(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_cbrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_cbrt(double x) {
   return cbrt(x);
 }
@@ -483,6 +571,11 @@ extern "C" __device__ double test_cbrt(double x) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.ceil.f32(float [[X:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test_ceilf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.ceil.f32(float [[X:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test_ceilf(float x) {
   return ceilf(x);
 }
@@ -497,6 +590,11 @@ extern "C" __device__ float test_ceilf(float x) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.ceil.f64(double [[X:%.*]])
 // FINITEONLY-NEXT:    ret double [[TMP0]]
 //
+// APPROX-LABEL: @test_ceil(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.ceil.f64(double [[X:%.*]])
+// APPROX-NEXT:    ret double [[TMP0]]
+//
 extern "C" __device__ double test_ceil(double x) {
   return ceil(x);
 }
@@ -511,6 +609,11 @@ extern "C" __device__ double test_ceil(double x) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.copysign.f32(float [[X:%.*]], float [[Y:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test_copysignf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.copysign.f32(float [[X:%.*]], float [[Y:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test_copysignf(float x, float y) {
   return copysignf(x, y);
 }
@@ -525,6 +628,11 @@ extern "C" __device__ float test_copysignf(float x, float y) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.copysign.f64(double [[X:%.*]], double [[Y:%.*]])
 // FINITEONLY-NEXT:    ret double [[TMP0]]
 //
+// APPROX-LABEL: @test_copysign(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.copysign.f64(double [[X:%.*]], double [[Y:%.*]])
+// APPROX-NEXT:    ret double [[TMP0]]
+//
 extern "C" __device__ double test_copysign(double x, double y) {
   return copysign(x, y);
 }
@@ -539,6 +647,11 @@ extern "C" __device__ double test_copysign(double x, double y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_cos_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16:[0-9]+]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_cosf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR16:[0-9]+]]
+// APPROX-NEXT:    ret float [[CALL_I_I]]
+//
 extern "C" __device__ float test_cosf(float x) {
   return cosf(x);
 }
@@ -553,6 +666,11 @@ extern "C" __device__ float test_cosf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_cos_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_cos(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_cos_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_cos(double x) {
   return cos(x);
 }
@@ -567,6 +685,11 @@ extern "C" __device__ double test_cos(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_cosh_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_coshf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_cosh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_coshf(float x) {
   return coshf(x);
 }
@@ -581,6 +704,11 @@ extern "C" __device__ float test_coshf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_cosh_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_cosh(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_cosh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_cosh(double x) {
   return cosh(x);
 }
@@ -595,6 +723,11 @@ extern "C" __device__ double test_cosh(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_cospi_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_cospif(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_cospi_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_cospif(float x) {
   return cospif(x);
 }
@@ -609,6 +742,11 @@ extern "C" __device__ float test_cospif(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_cospi_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_cospi(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_cospi_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_cospi(double x) {
   return cospi(x);
 }
@@ -623,6 +761,11 @@ extern "C" __device__ double test_cospi(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_i0_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_cyl_bessel_i0f(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_i0_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_cyl_bessel_i0f(float x) {
   return cyl_bessel_i0f(x);
 }
@@ -637,6 +780,11 @@ extern "C" __device__ float test_cyl_bessel_i0f(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_i0_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_cyl_bessel_i0(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_i0_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_cyl_bessel_i0(double x) {
   return cyl_bessel_i0(x);
 }
@@ -651,6 +799,11 @@ extern "C" __device__ double test_cyl_bessel_i0(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_i1_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_cyl_bessel_i1f(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_i1_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_cyl_bessel_i1f(float x) {
   return cyl_bessel_i1f(x);
 }
@@ -665,6 +818,11 @@ extern "C" __device__ float test_cyl_bessel_i1f(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_i1_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_cyl_bessel_i1(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_i1_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_cyl_bessel_i1(double x) {
   return cyl_bessel_i1(x);
 }
@@ -679,6 +837,11 @@ extern "C" __device__ double test_cyl_bessel_i1(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_erfc_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_erfcf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_erfc_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_erfcf(float x) {
   return erfcf(x);
 }
@@ -693,6 +856,11 @@ extern "C" __device__ float test_erfcf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_erfc_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_erfc(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_erfc_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_erfc(double x) {
   return erfc(x);
 }
@@ -707,6 +875,11 @@ extern "C" __device__ double test_erfc(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_erfinv_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_erfinvf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_erfinv_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_erfinvf(float x) {
   return erfinvf(x);
 }
@@ -721,6 +894,11 @@ extern "C" __device__ float test_erfinvf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_erfinv_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_erfinv(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_erfinv_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_erfinv(double x) {
   return erfinv(x);
 }
@@ -735,6 +913,11 @@ extern "C" __device__ double test_erfinv(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_exp10_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_exp10f(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_exp10_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_exp10f(float x) {
   return exp10f(x);
 }
@@ -749,6 +932,11 @@ extern "C" __device__ float test_exp10f(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_exp10_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_exp10(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_exp10_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_exp10(double x) {
   return exp10(x);
 }
@@ -763,6 +951,11 @@ extern "C" __device__ double test_exp10(double x) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.exp2.f32(float [[X:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test_exp2f(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.exp2.f32(float [[X:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test_exp2f(float x) {
   return exp2f(x);
 }
@@ -777,6 +970,11 @@ extern "C" __device__ float test_exp2f(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_exp2_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_exp2(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_exp2_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_exp2(double x) {
   return exp2(x);
 }
@@ -791,6 +989,11 @@ extern "C" __device__ double test_exp2(double x) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.exp.f32(float [[X:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test_expf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.exp.f32(float [[X:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test_expf(float x) {
   return expf(x);
 }
@@ -805,6 +1008,11 @@ extern "C" __device__ float test_expf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_exp_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_exp(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_exp_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_exp(double x) {
   return exp(x);
 }
@@ -819,6 +1027,11 @@ extern "C" __device__ double test_exp(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_expm1_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_expm1f(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_expm1_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_expm1f(float x) {
   return expm1f(x);
 }
@@ -833,6 +1046,11 @@ extern "C" __device__ float test_expm1f(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_expm1_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_expm1(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_expm1_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_expm1(double x) {
   return expm1(x);
 }
@@ -847,6 +1065,11 @@ extern "C" __device__ double test_expm1(double x) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.fabs.f32(float [[X:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test_fabsf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.fabs.f32(float [[X:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test_fabsf(float x) {
   return fabsf(x);
 }
@@ -861,6 +1084,11 @@ extern "C" __device__ float test_fabsf(float x) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.fabs.f64(double [[X:%.*]])
 // FINITEONLY-NEXT:    ret double [[TMP0]]
 //
+// APPROX-LABEL: @test_fabs(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.fabs.f64(double [[X:%.*]])
+// APPROX-NEXT:    ret double [[TMP0]]
+//
 extern "C" __device__ double test_fabs(double x) {
   return fabs(x);
 }
@@ -875,6 +1103,11 @@ extern "C" __device__ double test_fabs(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fdim_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_fdimf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_fdim_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_fdimf(float x, float y) {
   return fdimf(x, y);
 }
@@ -889,6 +1122,11 @@ extern "C" __device__ float test_fdimf(float x, float y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fdim_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_fdim(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_fdim_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_fdim(double x, double y) {
   return fdim(x, y);
 }
@@ -903,6 +1141,11 @@ extern "C" __device__ double test_fdim(double x, double y) {
 // FINITEONLY-NEXT:    [[DIV_I:%.*]] = fdiv nnan ninf contract float [[X:%.*]], [[Y:%.*]]
 // FINITEONLY-NEXT:    ret float [[DIV_I]]
 //
+// APPROX-LABEL: @test_fdividef(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[DIV_I:%.*]] = fdiv contract float [[X:%.*]], [[Y:%.*]]
+// APPROX-NEXT:    ret float [[DIV_I]]
+//
 extern "C" __device__ float test_fdividef(float x, float y) {
   return fdividef(x, y);
 }
@@ -917,6 +1160,11 @@ extern "C" __device__ float test_fdividef(float x, float y) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.floor.f32(float [[X:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test_floorf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.floor.f32(float [[X:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test_floorf(float x) {
   return floorf(x);
 }
@@ -931,6 +1179,11 @@ extern "C" __device__ float test_floorf(float x) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.floor.f64(double [[X:%.*]])
 // FINITEONLY-NEXT:    ret double [[TMP0]]
 //
+// APPROX-LABEL: @test_floor(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.floor.f64(double [[X:%.*]])
+// APPROX-NEXT:    ret double [[TMP0]]
+//
 extern "C" __device__ double test_floor(double x) {
   return floor(x);
 }
@@ -945,6 +1198,11 @@ extern "C" __device__ double test_floor(double x) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test_fmaf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test_fmaf(float x, float y, float z) {
   return fmaf(x, y, z);
 }
@@ -959,6 +1217,11 @@ extern "C" __device__ float test_fmaf(float x, float y, float z) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
 // FINITEONLY-NEXT:    ret double [[TMP0]]
 //
+// APPROX-LABEL: @test_fma(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
+// APPROX-NEXT:    ret double [[TMP0]]
+//
 extern "C" __device__ double test_fma(double x, double y, double z) {
   return fma(x, y, z);
 }
@@ -973,6 +1236,11 @@ extern "C" __device__ double test_fma(double x, double y, double z) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
 // FINITEONLY-NEXT:    ret double [[TMP0]]
 //
+// APPROX-LABEL: @test_fma_rn(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
+// APPROX-NEXT:    ret double [[TMP0]]
+//
 extern "C" __device__ double test_fma_rn(double x, double y, double z) {
   return __fma_rn(x, y, z);
 }
@@ -987,6 +1255,11 @@ extern "C" __device__ double test_fma_rn(double x, double y, double z) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test_fmaxf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test_fmaxf(float x, float y) {
   return fmaxf(x, y);
 }
@@ -1001,6 +1274,11 @@ extern "C" __device__ float test_fmaxf(float x, float y) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
 // FINITEONLY-NEXT:    ret double [[TMP0]]
 //
+// APPROX-LABEL: @test_fmax(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// APPROX-NEXT:    ret double [[TMP0]]
+//
 extern "C" __device__ double test_fmax(double x, double y) {
   return fmax(x, y);
 }
@@ -1015,6 +1293,11 @@ extern "C" __device__ double test_fmax(double x, double y) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test_fminf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test_fminf(float x, float y) {
   return fminf(x, y);
 }
@@ -1029,6 +1312,11 @@ extern "C" __device__ float test_fminf(float x, float y) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
 // FINITEONLY-NEXT:    ret double [[TMP0]]
 //
+// APPROX-LABEL: @test_fmin(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// APPROX-NEXT:    ret double [[TMP0]]
+//
 extern "C" __device__ double test_fmin(double x, double y) {
   return fmin(x, y);
 }
@@ -1043,6 +1331,11 @@ extern "C" __device__ double test_fmin(double x, double y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_fmod_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_fmodf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_fmod_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_fmodf(float x, float y) {
   return fmodf(x, y);
 }
@@ -1057,6 +1350,11 @@ extern "C" __device__ float test_fmodf(float x, float y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_fmod_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_fmod(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_fmod_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_fmod(double x, double y) {
   return fmod(x, y);
 }
@@ -1095,6 +1393,11 @@ extern "C" __device__ double test_frexp(double x, int* y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_hypot_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_hypotf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_hypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_hypotf(float x, float y) {
   return hypotf(x, y);
 }
@@ -1109,6 +1412,11 @@ extern "C" __device__ float test_hypotf(float x, float y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_hypot_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_hypot(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_hypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_hypot(double x, double y) {
   return hypot(x, y);
 }
@@ -1123,6 +1431,11 @@ extern "C" __device__ double test_hypot(double x, double y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret i32 [[CALL_I]]
 //
+// APPROX-LABEL: @test_ilogbf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret i32 [[CALL_I]]
+//
 extern "C" __device__ int test_ilogbf(float x) {
   return ilogbf(x);
 }
@@ -1137,6 +1450,11 @@ extern "C" __device__ int test_ilogbf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret i32 [[CALL_I]]
 //
+// APPROX-LABEL: @test_ilogb(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call i32 @__ocml_ilogb_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret i32 [[CALL_I]]
+//
 extern "C" __device__ int test_ilogb(double x) {
   return ilogb(x);
 }
@@ -1151,6 +1469,12 @@ extern "C" __device__ int test_ilogb(double x) {
 // FINITEONLY-NEXT:  entry:
 // FINITEONLY-NEXT:    ret i32 1
 //
+// APPROX-LABEL: @test___finitef(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call i1 @llvm.is.fpclass.f32(float [[X:%.*]], i32 504)
+// APPROX-NEXT:    [[CONV:%.*]] = zext i1 [[TMP0]] to i32
+// APPROX-NEXT:    ret i32 [[CONV]]
+//
 extern "C" __device__ BOOL_TYPE test___finitef(float x) {
   return __finitef(x);
 }
@@ -1165,6 +1489,12 @@ extern "C" __device__ BOOL_TYPE test___finitef(float x) {
 // FINITEONLY-NEXT:  entry:
 // FINITEONLY-NEXT:    ret i32 1
 //
+// APPROX-LABEL: @test___finite(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call i1 @llvm.is.fpclass.f64(double [[X:%.*]], i32 504)
+// APPROX-NEXT:    [[CONV:%.*]] = zext i1 [[TMP0]] to i32
+// APPROX-NEXT:    ret i32 [[CONV]]
+//
 extern "C" __device__ BOOL_TYPE test___finite(double x) {
   return __finite(x);
 }
@@ -1179,6 +1509,12 @@ extern "C" __device__ BOOL_TYPE test___finite(double x) {
 // FINITEONLY-NEXT:  entry:
 // FINITEONLY-NEXT:    ret i32 0
 //
+// APPROX-LABEL: @test___isinff(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call i1 @llvm.is.fpclass.f32(float [[X:%.*]], i32 516)
+// APPROX-NEXT:    [[CONV:%.*]] = zext i1 [[TMP0]] to i32
+// APPROX-NEXT:    ret i32 [[CONV]]
+//
 extern "C" __device__ BOOL_TYPE test___isinff(float x) {
   return __isinff(x);
 }
@@ -1193,6 +1529,12 @@ extern "C" __device__ BOOL_TYPE test___isinff(float x) {
 // FINITEONLY-NEXT:  entry:
 // FINITEONLY-NEXT:    ret i32 0
 //
+// APPROX-LABEL: @test___isinf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call i1 @llvm.is.fpclass.f64(double [[X:%.*]], i32 516)
+// APPROX-NEXT:    [[CONV:%.*]] = zext i1 [[TMP0]] to i32
+// APPROX-NEXT:    ret i32 [[CONV]]
+//
 extern "C" __device__ BOOL_TYPE test___isinf(double x) {
   return __isinf(x);
 }
@@ -1207,6 +1549,12 @@ extern "C" __device__ BOOL_TYPE test___isinf(double x) {
 // FINITEONLY-NEXT:  entry:
 // FINITEONLY-NEXT:    ret i32 0
 //
+// APPROX-LABEL: @test___isnanf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = fcmp uno float [[X:%.*]], 0.000000e+00
+// APPROX-NEXT:    [[CONV:%.*]] = zext i1 [[TMP0]] to i32
+// APPROX-NEXT:    ret i32 [[CONV]]
+//
 extern "C" __device__ BOOL_TYPE test___isnanf(float x) {
   return __isnanf(x);
 }
@@ -1221,6 +1569,12 @@ extern "C" __device__ BOOL_TYPE test___isnanf(float x) {
 // FINITEONLY-NEXT:  entry:
 // FINITEONLY-NEXT:    ret i32 0
 //
+// APPROX-LABEL: @test___isnan(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = fcmp uno double [[X:%.*]], 0.000000e+00
+// APPROX-NEXT:    [[CONV:%.*]] = zext i1 [[TMP0]] to i32
+// APPROX-NEXT:    ret i32 [[CONV]]
+//
 extern "C" __device__ BOOL_TYPE test___isnan(double x) {
   return __isnan(x);
 }
@@ -1235,6 +1589,11 @@ extern "C" __device__ BOOL_TYPE test___isnan(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_j0_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_j0f(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_j0f(float x) {
   return j0f(x);
 }
@@ -1249,6 +1608,11 @@ extern "C" __device__ float test_j0f(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_j0_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_j0(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_j0(double x) {
   return j0(x);
 }
@@ -1263,6 +1627,11 @@ extern "C" __device__ double test_j0(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_j1_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_j1f(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_j1f(float x) {
   return j1f(x);
 }
@@ -1277,6 +1646,11 @@ extern "C" __device__ float test_j1f(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_j1_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_j1(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_j1(double x) {
   return j1(x);
 }
@@ -1347,6 +1721,39 @@ extern "C" __device__ double test_j1(double x) {
 // 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]]
 //
+// APPROX-LABEL: @test_jnf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
+// APPROX-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
+// APPROX-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
+// APPROX-NEXT:    ]
+// APPROX:       if.then.i:
+// APPROX-NEXT:    [[CALL_I_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    br label [[_ZL3JNFIF_EXIT:%.*]]
+// APPROX:       if.then2.i:
+// APPROX-NEXT:    [[CALL_I20_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT:    br label [[_ZL3JNFIF_EXIT]]
+// APPROX:       if.end4.i:
+// APPROX-NEXT:    [[CALL_I21_I:%.*]] = tail call contract float @__ocml_j0_f32(float noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT:    [[CALL_I22_I:%.*]] = tail call contract float @__ocml_j1_f32(float noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// APPROX-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3JNFIF_EXIT]]
+// APPROX:       for.body.i:
+// APPROX-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// APPROX-NEXT:    [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
+// APPROX-NEXT:    [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// APPROX-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// APPROX-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
+// APPROX-NEXT:    [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
+// APPROX-NEXT:    [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]]
+// APPROX-NEXT:    [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]]
+// APPROX-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
+// APPROX-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// APPROX-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL3JNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP14:![0-9]+]]
+// APPROX:       _ZL3jnfif.exit:
+// APPROX-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]] ]
+// APPROX-NEXT:    ret float [[RETVAL_0_I]]
+//
 extern "C" __device__ float test_jnf(int x, float y) {
   return jnf(x, y);
 }
@@ -1417,6 +1824,39 @@ extern "C" __device__ float test_jnf(int x, float y) {
 // 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]]
 //
+// APPROX-LABEL: @test_jn(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
+// APPROX-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
+// APPROX-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
+// APPROX-NEXT:    ]
+// APPROX:       if.then.i:
+// APPROX-NEXT:    [[CALL_I_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    br label [[_ZL2JNID_EXIT:%.*]]
+// APPROX:       if.then2.i:
+// APPROX-NEXT:    [[CALL_I20_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT:    br label [[_ZL2JNID_EXIT]]
+// APPROX:       if.end4.i:
+// APPROX-NEXT:    [[CALL_I21_I:%.*]] = tail call contract double @__ocml_j0_f64(double noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT:    [[CALL_I22_I:%.*]] = tail call contract double @__ocml_j1_f64(double noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// APPROX-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2JNID_EXIT]]
+// APPROX:       for.body.i:
+// APPROX-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// APPROX-NEXT:    [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
+// APPROX-NEXT:    [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// APPROX-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// APPROX-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
+// APPROX-NEXT:    [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
+// APPROX-NEXT:    [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]]
+// APPROX-NEXT:    [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]]
+// APPROX-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
+// APPROX-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// APPROX-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL2JNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP15:![0-9]+]]
+// APPROX:       _ZL2jnid.exit:
+// APPROX-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]] ]
+// APPROX-NEXT:    ret double [[RETVAL_0_I]]
+//
 extern "C" __device__ double test_jn(int x, double y) {
   return jn(x, y);
 }
@@ -1431,6 +1871,11 @@ extern "C" __device__ double test_jn(int x, double y) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[Y:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test_ldexpf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[Y:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test_ldexpf(float x, int y) {
   return ldexpf(x, y);
 }
@@ -1445,6 +1890,11 @@ extern "C" __device__ float test_ldexpf(float x, int y) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[Y:%.*]])
 // FINITEONLY-NEXT:    ret double [[TMP0]]
 //
+// APPROX-LABEL: @test_ldexp(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[Y:%.*]])
+// APPROX-NEXT:    ret double [[TMP0]]
+//
 extern "C" __device__ double test_ldexp(double x, int y) {
   return ldexp(x, y);
 }
@@ -1459,6 +1909,11 @@ extern "C" __device__ double test_ldexp(double x, int y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_lgamma_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_lgammaf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_lgamma_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_lgammaf(float x) {
   return lgammaf(x);
 }
@@ -1473,6 +1928,11 @@ extern "C" __device__ float test_lgammaf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_lgamma_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_lgamma(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_lgamma_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_lgamma(double x) {
   return lgamma(x);
 }
@@ -1489,6 +1949,12 @@ extern "C" __device__ double test_lgamma(double x) {
 // FINITEONLY-NEXT:    [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
 // FINITEONLY-NEXT:    ret i64 [[CONV_I]]
 //
+// APPROX-LABEL: @test_llrintf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.rint.f32(float [[X:%.*]])
+// APPROX-NEXT:    [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
+// APPROX-NEXT:    ret i64 [[CONV_I]]
+//
 extern "C" __device__ long long int test_llrintf(float x) {
   return llrintf(x);
 }
@@ -1505,6 +1971,12 @@ extern "C" __device__ long long int test_llrintf(float x) {
 // FINITEONLY-NEXT:    [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
 // FINITEONLY-NEXT:    ret i64 [[CONV_I]]
 //
+// APPROX-LABEL: @test_llrint(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.rint.f64(double [[X:%.*]])
+// APPROX-NEXT:    [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
+// APPROX-NEXT:    ret i64 [[CONV_I]]
+//
 extern "C" __device__ long long int test_llrint(double x) {
   return llrint(x);
 }
@@ -1521,6 +1993,12 @@ extern "C" __device__ long long int test_llrint(double x) {
 // FINITEONLY-NEXT:    [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
 // FINITEONLY-NEXT:    ret i64 [[CONV_I]]
 //
+// APPROX-LABEL: @test_llroundf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.round.f32(float [[X:%.*]])
+// APPROX-NEXT:    [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
+// APPROX-NEXT:    ret i64 [[CONV_I]]
+//
 extern "C" __device__ long long int test_llroundf(float x) {
   return llroundf(x);
 }
@@ -1537,6 +2015,12 @@ extern "C" __device__ long long int test_llroundf(float x) {
 // FINITEONLY-NEXT:    [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
 // FINITEONLY-NEXT:    ret i64 [[CONV_I]]
 //
+// APPROX-LABEL: @test_llround(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.round.f64(double [[X:%.*]])
+// APPROX-NEXT:    [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
+// APPROX-NEXT:    ret i64 [[CONV_I]]
+//
 extern "C" __device__ long long int test_llround(double x) {
   return llround(x);
 }
@@ -1551,6 +2035,11 @@ extern "C" __device__ long long int test_llround(double x) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.log10.f32(float [[X:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test_log10f(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.log10.f32(float [[X:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test_log10f(float x) {
   return log10f(x);
 }
@@ -1565,6 +2054,11 @@ extern "C" __device__ float test_log10f(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_log10_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_log10(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_log10_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_log10(double x) {
   return log10(x);
 }
@@ -1579,6 +2073,11 @@ extern "C" __device__ double test_log10(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_log1p_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_log1pf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_log1p_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_log1pf(float x) {
   return log1pf(x);
 }
@@ -1593,19 +2092,29 @@ extern "C" __device__ float test_log1pf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_log1p_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_log1p(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_log1p_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-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:    [[TMP0:%.*]] = tail call contract float @llvm.log2.f32(float [[X:%.*]])
-// DEFAULT-NEXT:    ret float [[TMP0]]
+// 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:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.log2.f32(float [[X:%.*]])
-// FINITEONLY-NEXT:    ret float [[TMP0]]
+// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_log2_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+// APPROX-LABEL: @test_log2f(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.log.f32(float [[X:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
 //
 extern "C" __device__ float test_log2f(float x) {
   return log2f(x);
@@ -1621,6 +2130,11 @@ extern "C" __device__ float test_log2f(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_log2_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_log2(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_log2_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_log2(double x) {
   return log2(x);
 }
@@ -1635,6 +2149,11 @@ extern "C" __device__ double test_log2(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_logb_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_logbf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_logb_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_logbf(float x) {
   return logbf(x);
 }
@@ -1649,10 +2168,34 @@ extern "C" __device__ float test_logbf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_logb_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_logb(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_logb_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_logb(double x) {
   return logb(x);
 }
 
+// DEFAULT-LABEL: @test_logf(
+// DEFAULT-NEXT:  entry:
+// DEFAULT-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_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 nofpclass(nan inf) float @__ocml_log_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
+// FINITEONLY-NEXT:    ret float [[CALL_I]]
+//
+// APPROX-LABEL: @test_logf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.log.f32(float [[X:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
+extern "C" __device__ float test_logf(float x) {
+  return logf(x);
+}
+
 // DEFAULT-LABEL: @test_lrintf(
 // DEFAULT-NEXT:  entry:
 // DEFAULT-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.rint.f32(float [[X:%.*]])
@@ -1665,6 +2208,12 @@ extern "C" __device__ double test_logb(double x) {
 // FINITEONLY-NEXT:    [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
 // FINITEONLY-NEXT:    ret i64 [[CONV_I]]
 //
+// APPROX-LABEL: @test_lrintf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.rint.f32(float [[X:%.*]])
+// APPROX-NEXT:    [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
+// APPROX-NEXT:    ret i64 [[CONV_I]]
+//
 extern "C" __device__ long int test_lrintf(float x) {
   return lrintf(x);
 }
@@ -1681,6 +2230,12 @@ extern "C" __device__ long int test_lrintf(float x) {
 // FINITEONLY-NEXT:    [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
 // FINITEONLY-NEXT:    ret i64 [[CONV_I]]
 //
+// APPROX-LABEL: @test_lrint(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.rint.f64(double [[X:%.*]])
+// APPROX-NEXT:    [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
+// APPROX-NEXT:    ret i64 [[CONV_I]]
+//
 extern "C" __device__ long int test_lrint(double x) {
   return lrint(x);
 }
@@ -1697,6 +2252,12 @@ extern "C" __device__ long int test_lrint(double x) {
 // FINITEONLY-NEXT:    [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
 // FINITEONLY-NEXT:    ret i64 [[CONV_I]]
 //
+// APPROX-LABEL: @test_lroundf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.round.f32(float [[X:%.*]])
+// APPROX-NEXT:    [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
+// APPROX-NEXT:    ret i64 [[CONV_I]]
+//
 extern "C" __device__ long int test_lroundf(float x) {
   return lroundf(x);
 }
@@ -1713,6 +2274,12 @@ extern "C" __device__ long int test_lroundf(float x) {
 // FINITEONLY-NEXT:    [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
 // FINITEONLY-NEXT:    ret i64 [[CONV_I]]
 //
+// APPROX-LABEL: @test_lround(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.round.f64(double [[X:%.*]])
+// APPROX-NEXT:    [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
+// APPROX-NEXT:    ret i64 [[CONV_I]]
+//
 extern "C" __device__ long int test_lround(double x) {
   return lround(x);
 }
@@ -1737,6 +2304,16 @@ extern "C" __device__ long int test_lround(double x) {
 // FINITEONLY-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_modff(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
+// APPROX-NEXT:    call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17:[0-9]+]]
+// APPROX-NEXT:    [[CALL_I:%.*]] = call contract float @__ocml_modf_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// APPROX-NEXT:    [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA16:![0-9]+]]
+// APPROX-NEXT:    store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA16]]
+// APPROX-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_modff(float x, float* y) {
   return modff(x, y);
 }
@@ -1761,6 +2338,16 @@ extern "C" __device__ float test_modff(float x, float* y) {
 // FINITEONLY-NEXT:    call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_modf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
+// APPROX-NEXT:    call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// APPROX-NEXT:    [[CALL_I:%.*]] = call contract double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// APPROX-NEXT:    [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA18:![0-9]+]]
+// APPROX-NEXT:    store double [[TMP0]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA18]]
+// APPROX-NEXT:    call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_modf(double x, double* y) {
   return modf(x, y);
 }
@@ -2004,6 +2591,11 @@ extern "C" __device__ double test_nan_fill() {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.nearbyint.f32(float [[X:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test_nearbyintf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.nearbyint.f32(float [[X:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test_nearbyintf(float x) {
   return nearbyintf(x);
 }
@@ -2018,6 +2610,11 @@ extern "C" __device__ float test_nearbyintf(float x) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.nearbyint.f64(double [[X:%.*]])
 // FINITEONLY-NEXT:    ret double [[TMP0]]
 //
+// APPROX-LABEL: @test_nearbyint(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.nearbyint.f64(double [[X:%.*]])
+// APPROX-NEXT:    ret double [[TMP0]]
+//
 extern "C" __device__ double test_nearbyint(double x) {
   return nearbyint(x);
 }
@@ -2032,6 +2629,11 @@ extern "C" __device__ double test_nearbyint(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_nextafter_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_nextafterf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_nextafter_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_nextafterf(float x, float y) {
   return nextafterf(x, y);
 }
@@ -2046,6 +2648,11 @@ extern "C" __device__ float test_nextafterf(float x, float y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_nextafter_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_nextafter(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_nextafter_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_nextafter(double x, double y) {
   return nextafter(x, y);
 }
@@ -2060,6 +2667,11 @@ extern "C" __device__ double test_nextafter(double x, double y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_len3_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], float noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_norm3df(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_len3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_norm3df(float x, float y, float z) {
   return norm3df(x, y, z);
 }
@@ -2074,6 +2686,11 @@ extern "C" __device__ float test_norm3df(float x, float y, float z) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_len3_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_norm3d(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_len3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_norm3d(double x, double y, double z) {
   return norm3d(x, y, z);
 }
@@ -2088,6 +2705,11 @@ extern "C" __device__ double test_norm3d(double x, double y, double z) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_len4_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], float noundef nofpclass(nan inf) [[Z:%.*]], float noundef nofpclass(nan inf) [[W:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_norm4df(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_len4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR14]]
+// APPROX-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);
 }
@@ -2102,6 +2724,11 @@ extern "C" __device__ float test_norm4df(float x, float y, float z, float w) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_len4_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]], double noundef nofpclass(nan inf) [[W:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_norm4d(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_len4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR14]]
+// APPROX-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);
 }
@@ -2116,6 +2743,11 @@ extern "C" __device__ double test_norm4d(double x, double y, double z, double w)
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_ncdf_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_normcdff(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_ncdf_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_normcdff(float x) {
   return normcdff(x);
 }
@@ -2130,6 +2762,11 @@ extern "C" __device__ float test_normcdff(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_ncdf_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_normcdf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_ncdf_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_normcdf(double x) {
   return normcdf(x);
 }
@@ -2144,6 +2781,11 @@ extern "C" __device__ double test_normcdf(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_ncdfinv_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_normcdfinvf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_ncdfinv_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_normcdfinvf(float x) {
   return normcdfinvf(x);
 }
@@ -2158,6 +2800,11 @@ extern "C" __device__ float test_normcdfinvf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_ncdfinv_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_normcdfinv(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_ncdfinv_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_normcdfinv(double x) {
   return normcdfinv(x);
 }
@@ -2202,6 +2849,26 @@ extern "C" __device__ double test_normcdfinv(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_sqrt_f32(float noundef nofpclass(nan inf) [[__R_0_LCSSA_I]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_normf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// APPROX-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL5NORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// APPROX:       while.body.i:
+// APPROX-NEXT:    [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// APPROX-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// APPROX-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// APPROX-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
+// APPROX-NEXT:    [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA16]]
+// APPROX-NEXT:    [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]]
+// APPROX-NEXT:    [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]]
+// APPROX-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// APPROX-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
+// APPROX-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL5NORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
+// APPROX:       _ZL5normfiPKf.exit:
+// APPROX-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_normf(int x, const float *y) {
   return normf(x, y);
 }
@@ -2246,6 +2913,26 @@ extern "C" __device__ float test_normf(int x, const float *y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_sqrt_f64(double noundef nofpclass(nan inf) [[__R_0_LCSSA_I]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_norm(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// APPROX-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL4NORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// APPROX:       while.body.i:
+// APPROX-NEXT:    [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// APPROX-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// APPROX-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// APPROX-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
+// APPROX-NEXT:    [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA18]]
+// APPROX-NEXT:    [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]]
+// APPROX-NEXT:    [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]]
+// APPROX-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// APPROX-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
+// APPROX-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL4NORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
+// APPROX:       _ZL4normiPKd.exit:
+// APPROX-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR14]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_norm(int x, const double *y) {
   return norm(x, y);
 }
@@ -2260,6 +2947,11 @@ extern "C" __device__ double test_norm(int x, const double *y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_pow_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_powf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_powf(float x, float y) {
   return powf(x, y);
 }
@@ -2274,6 +2966,11 @@ extern "C" __device__ float test_powf(float x, float y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_pow_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_pow(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_pow_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_pow(double x, double y) {
   return pow(x, y);
 }
@@ -2288,6 +2985,11 @@ extern "C" __device__ double test_pow(double x, double y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_pown_f32(float noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_powif(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_pown_f32(float noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_powif(float x, int y) {
   return powif(x, y);
 }
@@ -2302,6 +3004,11 @@ extern "C" __device__ float test_powif(float x, int y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_pown_f64(double noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_powi(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_pown_f64(double noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_powi(double x, int y) {
   return powi(x, y);
 }
@@ -2316,6 +3023,11 @@ extern "C" __device__ double test_powi(double x, int y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rcbrt_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_rcbrtf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_rcbrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_rcbrtf(float x) {
   return rcbrtf(x);
 }
@@ -2330,6 +3042,11 @@ extern "C" __device__ float test_rcbrtf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rcbrt_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_rcbrt(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_rcbrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_rcbrt(double x) {
   return rcbrt(x);
 }
@@ -2344,6 +3061,11 @@ extern "C" __device__ double test_rcbrt(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_remainder_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_remainderf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_remainder_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_remainderf(float x, float y) {
   return remainderf(x, y);
 }
@@ -2358,6 +3080,11 @@ extern "C" __device__ float test_remainderf(float x, float y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_remainder_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_remainder(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_remainder_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_remainder(double x, double y) {
   return remainder(x, y);
 }
@@ -2382,6 +3109,16 @@ extern "C" __device__ double test_remainder(double x, double y) {
 // FINITEONLY-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_remquof(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
+// APPROX-NEXT:    call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// APPROX-NEXT:    [[CALL_I:%.*]] = call contract float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// APPROX-NEXT:    [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA12]]
+// APPROX-NEXT:    store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA12]]
+// APPROX-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_remquof(float x, float y, int* z) {
   return remquof(x, y, z);
 }
@@ -2406,6 +3143,16 @@ extern "C" __device__ float test_remquof(float x, float y, int* z) {
 // FINITEONLY-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_remquo(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
+// APPROX-NEXT:    call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// APPROX-NEXT:    [[CALL_I:%.*]] = call contract double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// APPROX-NEXT:    [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA12]]
+// APPROX-NEXT:    store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA12]]
+// APPROX-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_remquo(double x, double y, int* z) {
   return remquo(x, y, z);
 }
@@ -2420,6 +3167,11 @@ extern "C" __device__ double test_remquo(double x, double y, int* z) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rhypot_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_rhypotf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_rhypot_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_rhypotf(float x, float y) {
   return rhypotf(x, y);
 }
@@ -2434,6 +3186,11 @@ extern "C" __device__ float test_rhypotf(float x, float y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rhypot_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_rhypot(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_rhypot_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_rhypot(double x, double y) {
   return rhypot(x, y);
 }
@@ -2448,6 +3205,11 @@ extern "C" __device__ double test_rhypot(double x, double y) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.rint.f32(float [[X:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test_rintf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.rint.f32(float [[X:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test_rintf(float x) {
   return rintf(x);
 }
@@ -2462,6 +3224,11 @@ extern "C" __device__ float test_rintf(float x) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.rint.f64(double [[X:%.*]])
 // FINITEONLY-NEXT:    ret double [[TMP0]]
 //
+// APPROX-LABEL: @test_rint(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.rint.f64(double [[X:%.*]])
+// APPROX-NEXT:    ret double [[TMP0]]
+//
 extern "C" __device__ double test_rint(double x) {
   return rint(x);
 }
@@ -2506,6 +3273,26 @@ extern "C" __device__ double test_rint(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rsqrt_f32(float noundef nofpclass(nan inf) [[__R_0_LCSSA_I]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_rnormf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// APPROX-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL6RNORMFIPKF_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// APPROX:       while.body.i:
+// APPROX-NEXT:    [[__R_08_I:%.*]] = phi float [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// APPROX-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// APPROX-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// APPROX-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
+// APPROX-NEXT:    [[TMP0:%.*]] = load float, ptr [[__A_ADDR_07_I]], align 4, !tbaa [[TBAA16]]
+// APPROX-NEXT:    [[MUL_I:%.*]] = fmul contract float [[TMP0]], [[TMP0]]
+// APPROX-NEXT:    [[ADD_I]] = fadd contract float [[__R_08_I]], [[MUL_I]]
+// APPROX-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds float, ptr [[__A_ADDR_07_I]], i64 1
+// APPROX-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
+// APPROX-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL6RNORMFIPKF_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP22:![0-9]+]]
+// APPROX:       _ZL6rnormfiPKf.exit:
+// APPROX-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi float [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[__R_0_LCSSA_I]]) #[[ATTR15]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_rnormf(int x, const float* y) {
   return rnormf(x, y);
 }
@@ -2550,6 +3337,26 @@ extern "C" __device__ float test_rnormf(int x, const float* y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rsqrt_f64(double noundef nofpclass(nan inf) [[__R_0_LCSSA_I]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_rnorm(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TOBOOL_NOT5_I:%.*]] = icmp eq i32 [[X:%.*]], 0
+// APPROX-NEXT:    br i1 [[TOBOOL_NOT5_I]], label [[_ZL5RNORMIPKD_EXIT:%.*]], label [[WHILE_BODY_I:%.*]]
+// APPROX:       while.body.i:
+// APPROX-NEXT:    [[__R_08_I:%.*]] = phi double [ [[ADD_I:%.*]], [[WHILE_BODY_I]] ], [ 0.000000e+00, [[ENTRY:%.*]] ]
+// APPROX-NEXT:    [[__A_ADDR_07_I:%.*]] = phi ptr [ [[INCDEC_PTR_I:%.*]], [[WHILE_BODY_I]] ], [ [[Y:%.*]], [[ENTRY]] ]
+// APPROX-NEXT:    [[__DIM_ADDR_06_I:%.*]] = phi i32 [ [[DEC_I:%.*]], [[WHILE_BODY_I]] ], [ [[X]], [[ENTRY]] ]
+// APPROX-NEXT:    [[DEC_I]] = add nsw i32 [[__DIM_ADDR_06_I]], -1
+// APPROX-NEXT:    [[TMP0:%.*]] = load double, ptr [[__A_ADDR_07_I]], align 8, !tbaa [[TBAA18]]
+// APPROX-NEXT:    [[MUL_I:%.*]] = fmul contract double [[TMP0]], [[TMP0]]
+// APPROX-NEXT:    [[ADD_I]] = fadd contract double [[__R_08_I]], [[MUL_I]]
+// APPROX-NEXT:    [[INCDEC_PTR_I]] = getelementptr inbounds double, ptr [[__A_ADDR_07_I]], i64 1
+// APPROX-NEXT:    [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
+// APPROX-NEXT:    br i1 [[TOBOOL_NOT_I]], label [[_ZL5RNORMIPKD_EXIT]], label [[WHILE_BODY_I]], !llvm.loop [[LOOP23:![0-9]+]]
+// APPROX:       _ZL5rnormiPKd.exit:
+// APPROX-NEXT:    [[__R_0_LCSSA_I:%.*]] = phi double [ 0.000000e+00, [[ENTRY]] ], [ [[ADD_I]], [[WHILE_BODY_I]] ]
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[__R_0_LCSSA_I]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_rnorm(int x, const double* y) {
   return rnorm(x, y);
 }
@@ -2564,6 +3371,11 @@ extern "C" __device__ double test_rnorm(int x, const double* y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rlen3_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], float noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_rnorm3df(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_rlen3_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_rnorm3df(float x, float y, float z) {
   return rnorm3df(x, y, z);
 }
@@ -2578,6 +3390,11 @@ extern "C" __device__ float test_rnorm3df(float x, float y, float z) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rlen3_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_rnorm3d(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_rlen3_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_rnorm3d(double x, double y, double z) {
   return rnorm3d(x, y, z);
 }
@@ -2592,6 +3409,11 @@ extern "C" __device__ double test_rnorm3d(double x, double y, double z) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rlen4_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], float noundef nofpclass(nan inf) [[Z:%.*]], float noundef nofpclass(nan inf) [[W:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_rnorm4df(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_rlen4_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], float noundef [[Z:%.*]], float noundef [[W:%.*]]) #[[ATTR14]]
+// APPROX-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);
 }
@@ -2606,6 +3428,11 @@ extern "C" __device__ float test_rnorm4df(float x, float y, float z, float w) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rlen4_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]], double noundef nofpclass(nan inf) [[W:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_rnorm4d(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_rlen4_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], double noundef [[Z:%.*]], double noundef [[W:%.*]]) #[[ATTR14]]
+// APPROX-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);
 }
@@ -2620,6 +3447,11 @@ extern "C" __device__ double test_rnorm4d(double x, double y, double z, double w
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.round.f32(float [[X:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test_roundf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.round.f32(float [[X:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test_roundf(float x) {
   return roundf(x);
 }
@@ -2634,6 +3466,11 @@ extern "C" __device__ float test_roundf(float x) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.round.f64(double [[X:%.*]])
 // FINITEONLY-NEXT:    ret double [[TMP0]]
 //
+// APPROX-LABEL: @test_round(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.round.f64(double [[X:%.*]])
+// APPROX-NEXT:    ret double [[TMP0]]
+//
 extern "C" __device__ double test_round(double x) {
   return round(x);
 }
@@ -2648,6 +3485,11 @@ extern "C" __device__ double test_round(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_rsqrt_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_rsqrtf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_rsqrtf(float x) {
   return rsqrtf(x);
 }
@@ -2662,6 +3504,11 @@ extern "C" __device__ float test_rsqrtf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_rsqrt_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_rsqrt(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_rsqrt_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_rsqrt(double x) {
   return rsqrt(x);
 }
@@ -2696,6 +3543,21 @@ extern "C" __device__ double test_rsqrt(double x) {
 // FINITEONLY-NEXT:    [[COND_I:%.*]] = phi nnan ninf contract float [ [[TMP0]], [[COND_TRUE_I]] ], [ [[CALL_I]], [[COND_FALSE_I]] ]
 // FINITEONLY-NEXT:    ret float [[COND_I]]
 //
+// APPROX-LABEL: @test_scalblnf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CMP_NOT_I:%.*]] = icmp eq i64 [[Y:%.*]], 9223372036854775807
+// APPROX-NEXT:    br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]]
+// APPROX:       cond.true.i:
+// APPROX-NEXT:    [[CONV_I:%.*]] = trunc i64 [[Y]] to i32
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[CONV_I]])
+// APPROX-NEXT:    br label [[_ZL8SCALBLNFFL_EXIT:%.*]]
+// APPROX:       cond.false.i:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_scalb_f32(float noundef [[X]], float noundef 0x43E0000000000000) #[[ATTR14]]
+// APPROX-NEXT:    br label [[_ZL8SCALBLNFFL_EXIT]]
+// APPROX:       _ZL8scalblnffl.exit:
+// APPROX-NEXT:    [[COND_I:%.*]] = phi contract float [ [[TMP0]], [[COND_TRUE_I]] ], [ [[CALL_I]], [[COND_FALSE_I]] ]
+// APPROX-NEXT:    ret float [[COND_I]]
+//
 extern "C" __device__ float test_scalblnf(float x, long int y) {
   return scalblnf(x, y);
 }
@@ -2730,6 +3592,21 @@ extern "C" __device__ float test_scalblnf(float x, long int y) {
 // FINITEONLY-NEXT:    [[COND_I:%.*]] = phi nnan ninf contract double [ [[TMP0]], [[COND_TRUE_I]] ], [ [[CALL_I]], [[COND_FALSE_I]] ]
 // FINITEONLY-NEXT:    ret double [[COND_I]]
 //
+// APPROX-LABEL: @test_scalbln(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CMP_NOT_I:%.*]] = icmp eq i64 [[Y:%.*]], 9223372036854775807
+// APPROX-NEXT:    br i1 [[CMP_NOT_I]], label [[COND_FALSE_I:%.*]], label [[COND_TRUE_I:%.*]]
+// APPROX:       cond.true.i:
+// APPROX-NEXT:    [[CONV_I:%.*]] = trunc i64 [[Y]] to i32
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[CONV_I]])
+// APPROX-NEXT:    br label [[_ZL7SCALBLNDL_EXIT:%.*]]
+// APPROX:       cond.false.i:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_scalb_f64(double noundef [[X]], double noundef 0x43E0000000000000) #[[ATTR14]]
+// APPROX-NEXT:    br label [[_ZL7SCALBLNDL_EXIT]]
+// APPROX:       _ZL7scalblndl.exit:
+// APPROX-NEXT:    [[COND_I:%.*]] = phi contract double [ [[TMP0]], [[COND_TRUE_I]] ], [ [[CALL_I]], [[COND_FALSE_I]] ]
+// APPROX-NEXT:    ret double [[COND_I]]
+//
 extern "C" __device__ double test_scalbln(double x, long int y) {
   return scalbln(x, y);
 }
@@ -2744,6 +3621,11 @@ extern "C" __device__ double test_scalbln(double x, long int y) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[Y:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test_scalbnf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[Y:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test_scalbnf(float x, int y) {
   return scalbnf(x, y);
 }
@@ -2758,6 +3640,11 @@ extern "C" __device__ float test_scalbnf(float x, int y) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[Y:%.*]])
 // FINITEONLY-NEXT:    ret double [[TMP0]]
 //
+// APPROX-LABEL: @test_scalbn(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[Y:%.*]])
+// APPROX-NEXT:    ret double [[TMP0]]
+//
 extern "C" __device__ double test_scalbn(double x, int y) {
   return scalbn(x, y);
 }
@@ -2805,6 +3692,14 @@ extern "C" __device__ BOOL_TYPE test___signbit(double x) {
 // FINITEONLY-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
 // FINITEONLY-NEXT:    ret void
 //
+// APPROX-LABEL: @test_sincosf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    store float [[CALL_I_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA16]]
+// APPROX-NEXT:    [[CALL1_I_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]]
+// APPROX-NEXT:    store float [[CALL1_I_I]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA16]]
+// APPROX-NEXT:    ret void
+//
 extern "C" __device__ void test_sincosf(float x, float *y, float *z) {
   sincosf(x, y, z);
 }
@@ -2831,6 +3726,17 @@ extern "C" __device__ void test_sincosf(float x, float *y, float *z) {
 // FINITEONLY-NEXT:    call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
 // FINITEONLY-NEXT:    ret void
 //
+// APPROX-LABEL: @test_sincos(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
+// APPROX-NEXT:    call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// APPROX-NEXT:    [[CALL_I:%.*]] = call contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// APPROX-NEXT:    store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA18]]
+// APPROX-NEXT:    [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA18]]
+// APPROX-NEXT:    store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA18]]
+// APPROX-NEXT:    call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// APPROX-NEXT:    ret void
+//
 extern "C" __device__ void test_sincos(double x, double *y, double *z) {
   sincos(x, y, z);
 }
@@ -2857,6 +3763,17 @@ extern "C" __device__ void test_sincos(double x, double *y, double *z) {
 // FINITEONLY-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
 // FINITEONLY-NEXT:    ret void
 //
+// APPROX-LABEL: @test_sincospif(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
+// APPROX-NEXT:    call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// APPROX-NEXT:    [[CALL_I:%.*]] = call contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// APPROX-NEXT:    store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA16]]
+// APPROX-NEXT:    [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA16]]
+// APPROX-NEXT:    store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA16]]
+// APPROX-NEXT:    call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// APPROX-NEXT:    ret void
+//
 extern "C" __device__ void test_sincospif(float x, float *y, float *z) {
   sincospif(x, y, z);
 }
@@ -2883,6 +3800,17 @@ extern "C" __device__ void test_sincospif(float x, float *y, float *z) {
 // FINITEONLY-NEXT:    call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
 // FINITEONLY-NEXT:    ret void
 //
+// APPROX-LABEL: @test_sincospi(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
+// APPROX-NEXT:    call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// APPROX-NEXT:    [[CALL_I:%.*]] = call contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
+// APPROX-NEXT:    store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA18]]
+// APPROX-NEXT:    [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA18]]
+// APPROX-NEXT:    store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA18]]
+// APPROX-NEXT:    call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
+// APPROX-NEXT:    ret void
+//
 extern "C" __device__ void test_sincospi(double x, double *y, double *z) {
   sincospi(x, y, z);
 }
@@ -2897,6 +3825,11 @@ extern "C" __device__ void test_sincospi(double x, double *y, double *z) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_sin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_sinf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret float [[CALL_I_I]]
+//
 extern "C" __device__ float test_sinf(float x) {
   return sinf(x);
 }
@@ -2911,6 +3844,11 @@ extern "C" __device__ float test_sinf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_sin_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_sin(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sin_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_sin(double x) {
   return sin(x);
 }
@@ -2925,6 +3863,11 @@ extern "C" __device__ double test_sin(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_sinpi_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_sinpif(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sinpi_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_sinpif(float x) {
   return sinpif(x);
 }
@@ -2939,6 +3882,11 @@ extern "C" __device__ float test_sinpif(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_sinpi_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_sinpi(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sinpi_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_sinpi(double x) {
   return sinpi(x);
 }
@@ -2953,6 +3901,11 @@ extern "C" __device__ double test_sinpi(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_sqrt_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_sqrtf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_sqrtf(float x) {
   return sqrtf(x);
 }
@@ -2967,6 +3920,11 @@ extern "C" __device__ float test_sqrtf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_sqrt_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_sqrt(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_sqrt(double x) {
   return sqrt(x);
 }
@@ -2981,6 +3939,11 @@ extern "C" __device__ double test_sqrt(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_tan_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_tanf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_tan_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_tanf(float x) {
   return tanf(x);
 }
@@ -2995,6 +3958,11 @@ extern "C" __device__ float test_tanf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_tan_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_tan(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_tan_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_tan(double x) {
   return tan(x);
 }
@@ -3009,6 +3977,11 @@ extern "C" __device__ double test_tan(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_tanh_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_tanhf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_tanh_f32(float noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_tanhf(float x) {
   return tanhf(x);
 }
@@ -3023,6 +3996,11 @@ extern "C" __device__ float test_tanhf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_tanh_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_tanh(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_tanh_f64(double noundef [[X:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_tanh(double x) {
   return tanh(x);
 }
@@ -3037,6 +4015,11 @@ extern "C" __device__ double test_tanh(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_tgamma_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_tgammaf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_tgamma_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_tgammaf(float x) {
   return tgammaf(x);
 }
@@ -3051,6 +4034,11 @@ extern "C" __device__ float test_tgammaf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_tgamma_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_tgamma(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_tgamma_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_tgamma(double x) {
   return tgamma(x);
 }
@@ -3065,6 +4053,11 @@ extern "C" __device__ double test_tgamma(double x) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.trunc.f32(float [[X:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test_truncf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.trunc.f32(float [[X:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test_truncf(float x) {
   return truncf(x);
 }
@@ -3079,6 +4072,11 @@ extern "C" __device__ float test_truncf(float x) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.trunc.f64(double [[X:%.*]])
 // FINITEONLY-NEXT:    ret double [[TMP0]]
 //
+// APPROX-LABEL: @test_trunc(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.trunc.f64(double [[X:%.*]])
+// APPROX-NEXT:    ret double [[TMP0]]
+//
 extern "C" __device__ double test_trunc(double x) {
   return trunc(x);
 }
@@ -3093,6 +4091,11 @@ extern "C" __device__ double test_trunc(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_y0_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_y0f(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_y0f(float x) {
   return y0f(x);
 }
@@ -3107,6 +4110,11 @@ extern "C" __device__ float test_y0f(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_y0_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_y0(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_y0(double x) {
   return y0(x);
 }
@@ -3121,6 +4129,11 @@ extern "C" __device__ double test_y0(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_y1_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test_y1f(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test_y1f(float x) {
   return y1f(x);
 }
@@ -3135,6 +4148,11 @@ extern "C" __device__ float test_y1f(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_y1_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test_y1(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test_y1(double x) {
   return y1(x);
 }
@@ -3205,6 +4223,39 @@ extern "C" __device__ double test_y1(double x) {
 // 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]]
 //
+// APPROX-LABEL: @test_ynf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
+// APPROX-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
+// APPROX-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
+// APPROX-NEXT:    ]
+// APPROX:       if.then.i:
+// APPROX-NEXT:    [[CALL_I_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    br label [[_ZL3YNFIF_EXIT:%.*]]
+// APPROX:       if.then2.i:
+// APPROX-NEXT:    [[CALL_I20_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT:    br label [[_ZL3YNFIF_EXIT]]
+// APPROX:       if.end4.i:
+// APPROX-NEXT:    [[CALL_I21_I:%.*]] = tail call contract float @__ocml_y0_f32(float noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT:    [[CALL_I22_I:%.*]] = tail call contract float @__ocml_y1_f32(float noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// APPROX-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL3YNFIF_EXIT]]
+// APPROX:       for.body.i:
+// APPROX-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// APPROX-NEXT:    [[__X1_025_I:%.*]] = phi float [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
+// APPROX-NEXT:    [[__X0_024_I:%.*]] = phi float [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// APPROX-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// APPROX-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to float
+// APPROX-NEXT:    [[DIV_I:%.*]] = fdiv contract float [[CONV_I]], [[Y]]
+// APPROX-NEXT:    [[MUL8_I:%.*]] = fmul contract float [[__X1_025_I]], [[DIV_I]]
+// APPROX-NEXT:    [[SUB_I]] = fsub contract float [[MUL8_I]], [[__X0_024_I]]
+// APPROX-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
+// APPROX-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// APPROX-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL3YNFIF_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP24:![0-9]+]]
+// APPROX:       _ZL3ynfif.exit:
+// APPROX-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]] ]
+// APPROX-NEXT:    ret float [[RETVAL_0_I]]
+//
 extern "C" __device__ float test_ynf(int x, float y) {
   return ynf(x, y);
 }
@@ -3275,6 +4326,39 @@ extern "C" __device__ float test_ynf(int x, float y) {
 // 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]]
 //
+// APPROX-LABEL: @test_yn(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    switch i32 [[X:%.*]], label [[IF_END4_I:%.*]] [
+// APPROX-NEXT:    i32 0, label [[IF_THEN_I:%.*]]
+// APPROX-NEXT:    i32 1, label [[IF_THEN2_I:%.*]]
+// APPROX-NEXT:    ]
+// APPROX:       if.then.i:
+// APPROX-NEXT:    [[CALL_I_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    br label [[_ZL2YNID_EXIT:%.*]]
+// APPROX:       if.then2.i:
+// APPROX-NEXT:    [[CALL_I20_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT:    br label [[_ZL2YNID_EXIT]]
+// APPROX:       if.end4.i:
+// APPROX-NEXT:    [[CALL_I21_I:%.*]] = tail call contract double @__ocml_y0_f64(double noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT:    [[CALL_I22_I:%.*]] = tail call contract double @__ocml_y1_f64(double noundef [[Y]]) #[[ATTR16]]
+// APPROX-NEXT:    [[CMP723_I:%.*]] = icmp sgt i32 [[X]], 1
+// APPROX-NEXT:    br i1 [[CMP723_I]], label [[FOR_BODY_I:%.*]], label [[_ZL2YNID_EXIT]]
+// APPROX:       for.body.i:
+// APPROX-NEXT:    [[__I_026_I:%.*]] = phi i32 [ [[INC_I:%.*]], [[FOR_BODY_I]] ], [ 1, [[IF_END4_I]] ]
+// APPROX-NEXT:    [[__X1_025_I:%.*]] = phi double [ [[SUB_I:%.*]], [[FOR_BODY_I]] ], [ [[CALL_I22_I]], [[IF_END4_I]] ]
+// APPROX-NEXT:    [[__X0_024_I:%.*]] = phi double [ [[__X1_025_I]], [[FOR_BODY_I]] ], [ [[CALL_I21_I]], [[IF_END4_I]] ]
+// APPROX-NEXT:    [[MUL_I:%.*]] = shl nuw nsw i32 [[__I_026_I]], 1
+// APPROX-NEXT:    [[CONV_I:%.*]] = sitofp i32 [[MUL_I]] to double
+// APPROX-NEXT:    [[DIV_I:%.*]] = fdiv contract double [[CONV_I]], [[Y]]
+// APPROX-NEXT:    [[MUL8_I:%.*]] = fmul contract double [[__X1_025_I]], [[DIV_I]]
+// APPROX-NEXT:    [[SUB_I]] = fsub contract double [[MUL8_I]], [[__X0_024_I]]
+// APPROX-NEXT:    [[INC_I]] = add nuw nsw i32 [[__I_026_I]], 1
+// APPROX-NEXT:    [[EXITCOND_NOT_I:%.*]] = icmp eq i32 [[INC_I]], [[X]]
+// APPROX-NEXT:    br i1 [[EXITCOND_NOT_I]], label [[_ZL2YNID_EXIT]], label [[FOR_BODY_I]], !llvm.loop [[LOOP25:![0-9]+]]
+// APPROX:       _ZL2ynid.exit:
+// APPROX-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]] ]
+// APPROX-NEXT:    ret double [[RETVAL_0_I]]
+//
 extern "C" __device__ double test_yn(int x, double y) {
   return yn(x, y);
 }
@@ -3289,19 +4373,32 @@ extern "C" __device__ double test_yn(int x, double y) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_cos_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test___cosf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-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]]
+// DEFAULT-NEXT:    [[MUL_I:%.*]] = fmul contract float [[X:%.*]], 0x400A934F00000000
+// DEFAULT-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.exp2.f32(float [[MUL_I]])
+// DEFAULT-NEXT:    ret float [[TMP0]]
 //
 // FINITEONLY-LABEL: @test___exp10f(
 // FINITEONLY-NEXT:  entry:
-// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_exp10_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
-// FINITEONLY-NEXT:    ret float [[CALL_I]]
+// FINITEONLY-NEXT:    [[MUL_I:%.*]] = fmul nnan ninf contract float [[X:%.*]], 0x400A934F00000000
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.amdgcn.exp2.f32(float [[MUL_I]])
+// FINITEONLY-NEXT:    ret float [[TMP0]]
+//
+// APPROX-LABEL: @test___exp10f(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[MUL_I:%.*]] = fmul contract float [[X:%.*]], 0x400A934F00000000
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.exp2.f32(float [[MUL_I]])
+// APPROX-NEXT:    ret float [[TMP0]]
 //
 extern "C" __device__ float test___exp10f(float x) {
   return __exp10f(x);
@@ -3309,13 +4406,21 @@ extern "C" __device__ float test___exp10f(float 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]]
+// DEFAULT-NEXT:    [[MUL_I:%.*]] = fmul contract float [[X:%.*]], 0x3FF7154760000000
+// DEFAULT-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.exp2.f32(float [[MUL_I]])
+// DEFAULT-NEXT:    ret float [[TMP0]]
 //
 // FINITEONLY-LABEL: @test___expf(
 // FINITEONLY-NEXT:  entry:
-// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_exp_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
-// FINITEONLY-NEXT:    ret float [[CALL_I]]
+// FINITEONLY-NEXT:    [[MUL_I:%.*]] = fmul nnan ninf contract float [[X:%.*]], 0x3FF7154760000000
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.amdgcn.exp2.f32(float [[MUL_I]])
+// FINITEONLY-NEXT:    ret float [[TMP0]]
+//
+// APPROX-LABEL: @test___expf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[MUL_I:%.*]] = fmul contract float [[X:%.*]], 0x3FF7154760000000
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.exp2.f32(float [[MUL_I]])
+// APPROX-NEXT:    ret float [[TMP0]]
 //
 extern "C" __device__ float test___expf(float x) {
   return __expf(x);
@@ -3331,6 +4436,11 @@ extern "C" __device__ float test___expf(float x) {
 // FINITEONLY-NEXT:    [[ADD_I:%.*]] = fadd nnan ninf contract float [[X:%.*]], [[Y:%.*]]
 // FINITEONLY-NEXT:    ret float [[ADD_I]]
 //
+// APPROX-LABEL: @test___fadd_rn(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[ADD_I:%.*]] = fadd contract float [[X:%.*]], [[Y:%.*]]
+// APPROX-NEXT:    ret float [[ADD_I]]
+//
 extern "C" __device__ float test___fadd_rn(float x, float y) {
   return __fadd_rn(x, y);
 }
@@ -3345,6 +4455,11 @@ extern "C" __device__ float test___fadd_rn(float x, float y) {
 // FINITEONLY-NEXT:    [[DIV_I:%.*]] = fdiv nnan ninf contract float [[X:%.*]], [[Y:%.*]]
 // FINITEONLY-NEXT:    ret float [[DIV_I]]
 //
+// APPROX-LABEL: @test___fdividef(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[DIV_I:%.*]] = fdiv contract float [[X:%.*]], [[Y:%.*]]
+// APPROX-NEXT:    ret float [[DIV_I]]
+//
 extern "C" __device__ float test___fdividef(float x, float y) {
   return __fdividef(x, y);
 }
@@ -3359,6 +4474,11 @@ extern "C" __device__ float test___fdividef(float x, float y) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test__fmaf_rn(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test__fmaf_rn(float x, float y, float z) {
   return __fmaf_rn(x, y, z);
 }
@@ -3373,6 +4493,11 @@ extern "C" __device__ float test__fmaf_rn(float x, float y, float z) {
 // FINITEONLY-NEXT:    [[MUL_I:%.*]] = fmul nnan ninf contract float [[X:%.*]], [[Y:%.*]]
 // FINITEONLY-NEXT:    ret float [[MUL_I]]
 //
+// APPROX-LABEL: @test___fmul_rn(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[MUL_I:%.*]] = fmul contract float [[X:%.*]], [[Y:%.*]]
+// APPROX-NEXT:    ret float [[MUL_I]]
+//
 extern "C" __device__ float test___fmul_rn(float x, float y) {
   return __fmul_rn(x, y);
 }
@@ -3387,6 +4512,11 @@ extern "C" __device__ float test___fmul_rn(float x, float y) {
 // FINITEONLY-NEXT:    [[DIV_I:%.*]] = fdiv nnan ninf contract float 1.000000e+00, [[X:%.*]]
 // FINITEONLY-NEXT:    ret float [[DIV_I]]
 //
+// APPROX-LABEL: @test___frcp_rn(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[DIV_I:%.*]] = fdiv contract float 1.000000e+00, [[X:%.*]]
+// APPROX-NEXT:    ret float [[DIV_I]]
+//
 extern "C" __device__ float test___frcp_rn(float x) {
   return __frcp_rn(x);
 }
@@ -3401,6 +4531,11 @@ extern "C" __device__ float test___frcp_rn(float x) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.amdgcn.rsq.f32(float [[X:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test___frsqrt_rn(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.rsq.f32(float [[X:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test___frsqrt_rn(float x) {
   return __frsqrt_rn(x);
 }
@@ -3415,6 +4550,11 @@ extern "C" __device__ float test___frsqrt_rn(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_sqrt_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test___fsqrt_rn(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_native_sqrt_f32(float noundef [[X:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test___fsqrt_rn(float x) {
   return __fsqrt_rn(x);
 }
@@ -3429,19 +4569,29 @@ extern "C" __device__ float test___fsqrt_rn(float x) {
 // FINITEONLY-NEXT:    [[SUB_I:%.*]] = fsub nnan ninf contract float [[X:%.*]], [[Y:%.*]]
 // FINITEONLY-NEXT:    ret float [[SUB_I]]
 //
+// APPROX-LABEL: @test___fsub_rn(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[SUB_I:%.*]] = fsub contract float [[X:%.*]], [[Y:%.*]]
+// APPROX-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]]
+// DEFAULT-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.log10.f32(float [[X:%.*]])
+// DEFAULT-NEXT:    ret float [[TMP0]]
 //
 // FINITEONLY-LABEL: @test___log10f(
 // FINITEONLY-NEXT:  entry:
-// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_log10_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
-// FINITEONLY-NEXT:    ret float [[CALL_I]]
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.log10.f32(float [[X:%.*]])
+// FINITEONLY-NEXT:    ret float [[TMP0]]
+//
+// APPROX-LABEL: @test___log10f(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.log10.f32(float [[X:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
 //
 extern "C" __device__ float test___log10f(float x) {
   return __log10f(x);
@@ -3449,13 +4599,18 @@ extern "C" __device__ float test___log10f(float 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]]
+// DEFAULT-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.log.f32(float [[X:%.*]])
+// DEFAULT-NEXT:    ret float [[TMP0]]
 //
 // FINITEONLY-LABEL: @test___log2f(
 // FINITEONLY-NEXT:  entry:
-// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_log2_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
-// FINITEONLY-NEXT:    ret float [[CALL_I]]
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.amdgcn.log.f32(float [[X:%.*]])
+// FINITEONLY-NEXT:    ret float [[TMP0]]
+//
+// APPROX-LABEL: @test___log2f(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.log.f32(float [[X:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
 //
 extern "C" __device__ float test___log2f(float x) {
   return __log2f(x);
@@ -3463,13 +4618,18 @@ extern "C" __device__ float test___log2f(float 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]]
+// DEFAULT-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.log.f32(float [[X:%.*]])
+// DEFAULT-NEXT:    ret float [[TMP0]]
 //
 // FINITEONLY-LABEL: @test___logf(
 // FINITEONLY-NEXT:  entry:
-// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_log_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR15]]
-// FINITEONLY-NEXT:    ret float [[CALL_I]]
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.log.f32(float [[X:%.*]])
+// FINITEONLY-NEXT:    ret float [[TMP0]]
+//
+// APPROX-LABEL: @test___logf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.log.f32(float [[X:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
 //
 extern "C" __device__ float test___logf(float x) {
   return __logf(x);
@@ -3485,6 +4645,11 @@ extern "C" __device__ float test___logf(float x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_pow_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR15]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test___powf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_pow_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR15]]
+// APPROX-NEXT:    ret float [[CALL_I]]
+//
 extern "C" __device__ float test___powf(float x, float y) {
   return __powf(x, y);
 }
@@ -3505,6 +4670,14 @@ extern "C" __device__ float test___powf(float x, float y) {
 // FINITEONLY-NEXT:    [[COND5_I:%.*]] = select nnan ninf contract i1 [[CMP_I]], float 0.000000e+00, float [[COND_I]]
 // FINITEONLY-NEXT:    ret float [[COND5_I]]
 //
+// APPROX-LABEL: @test___saturatef(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CMP_I:%.*]] = fcmp contract olt float [[X:%.*]], 0.000000e+00
+// APPROX-NEXT:    [[CMP1_I:%.*]] = fcmp contract ogt float [[X]], 1.000000e+00
+// APPROX-NEXT:    [[COND_I:%.*]] = select contract i1 [[CMP1_I]], float 1.000000e+00, float [[X]]
+// APPROX-NEXT:    [[COND5_I:%.*]] = select contract i1 [[CMP_I]], float 0.000000e+00, float [[COND_I]]
+// APPROX-NEXT:    ret float [[COND5_I]]
+//
 extern "C" __device__ float test___saturatef(float x) {
   return __saturatef(x);
 }
@@ -3525,6 +4698,14 @@ extern "C" __device__ float test___saturatef(float x) {
 // FINITEONLY-NEXT:    store float [[CALL1_I]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA16]]
 // FINITEONLY-NEXT:    ret void
 //
+// APPROX-LABEL: @test___sincosf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA16]]
+// APPROX-NEXT:    [[CALL1_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]]
+// APPROX-NEXT:    store float [[CALL1_I]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA16]]
+// APPROX-NEXT:    ret void
+//
 extern "C" __device__ void test___sincosf(float x, float *y, float *z) {
   __sincosf(x, y, z);
 }
@@ -3539,19 +4720,38 @@ extern "C" __device__ void test___sincosf(float x, float *y, float *z) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_sin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
 // FINITEONLY-NEXT:    ret float [[CALL_I]]
 //
+// APPROX-LABEL: @test___sinf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-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]]
+// DEFAULT-NEXT:    [[CALL_I_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// DEFAULT-NEXT:    [[CALL_I3_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]]
+// DEFAULT-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.rcp.f32(float [[CALL_I3_I]])
+// DEFAULT-NEXT:    [[MUL_I:%.*]] = fmul contract float [[CALL_I_I]], [[TMP0]]
+// DEFAULT-NEXT:    ret float [[MUL_I]]
 //
 // FINITEONLY-LABEL: @test___tanf(
 // FINITEONLY-NEXT:  entry:
-// FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_tan_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
-// FINITEONLY-NEXT:    ret float [[CALL_I]]
+// FINITEONLY-NEXT:    [[CALL_I_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_sin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    [[CALL_I3_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_native_cos_f32(float noundef nofpclass(nan inf) [[X]]) #[[ATTR16]]
+// FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.amdgcn.rcp.f32(float [[CALL_I3_I]])
+// FINITEONLY-NEXT:    [[MUL_I:%.*]] = fmul nnan ninf contract float [[CALL_I_I]], [[TMP0]]
+// FINITEONLY-NEXT:    ret float [[MUL_I]]
+//
+// APPROX-LABEL: @test___tanf(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I_I:%.*]] = tail call contract float @__ocml_native_sin_f32(float noundef [[X:%.*]]) #[[ATTR16]]
+// APPROX-NEXT:    [[CALL_I3_I:%.*]] = tail call contract float @__ocml_native_cos_f32(float noundef [[X]]) #[[ATTR16]]
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.amdgcn.rcp.f32(float [[CALL_I3_I]])
+// APPROX-NEXT:    [[MUL_I:%.*]] = fmul contract float [[CALL_I_I]], [[TMP0]]
+// APPROX-NEXT:    ret float [[MUL_I]]
 //
 extern "C" __device__ float test___tanf(float x) {
   return __tanf(x);
@@ -3567,6 +4767,11 @@ extern "C" __device__ float test___tanf(float x) {
 // FINITEONLY-NEXT:    [[ADD_I:%.*]] = fadd nnan ninf contract double [[X:%.*]], [[Y:%.*]]
 // FINITEONLY-NEXT:    ret double [[ADD_I]]
 //
+// APPROX-LABEL: @test___dadd_rn(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[ADD_I:%.*]] = fadd contract double [[X:%.*]], [[Y:%.*]]
+// APPROX-NEXT:    ret double [[ADD_I]]
+//
 extern "C" __device__ double test___dadd_rn(double x, double y) {
   return __dadd_rn(x, y);
 }
@@ -3581,6 +4786,11 @@ extern "C" __device__ double test___dadd_rn(double x, double y) {
 // FINITEONLY-NEXT:    [[DIV_I:%.*]] = fdiv nnan ninf contract double [[X:%.*]], [[Y:%.*]]
 // FINITEONLY-NEXT:    ret double [[DIV_I]]
 //
+// APPROX-LABEL: @test___ddiv_rn(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[DIV_I:%.*]] = fdiv contract double [[X:%.*]], [[Y:%.*]]
+// APPROX-NEXT:    ret double [[DIV_I]]
+//
 extern "C" __device__ double test___ddiv_rn(double x, double y) {
   return __ddiv_rn(x, y);
 }
@@ -3595,6 +4805,11 @@ extern "C" __device__ double test___ddiv_rn(double x, double y) {
 // FINITEONLY-NEXT:    [[MUL_I:%.*]] = fmul nnan ninf contract double [[X:%.*]], [[Y:%.*]]
 // FINITEONLY-NEXT:    ret double [[MUL_I]]
 //
+// APPROX-LABEL: @test___dmul_rn(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[MUL_I:%.*]] = fmul contract double [[X:%.*]], [[Y:%.*]]
+// APPROX-NEXT:    ret double [[MUL_I]]
+//
 extern "C" __device__ double test___dmul_rn(double x, double y) {
   return __dmul_rn(x, y);
 }
@@ -3609,6 +4824,11 @@ extern "C" __device__ double test___dmul_rn(double x, double y) {
 // FINITEONLY-NEXT:    [[DIV_I:%.*]] = fdiv nnan ninf contract double 1.000000e+00, [[X:%.*]]
 // FINITEONLY-NEXT:    ret double [[DIV_I]]
 //
+// APPROX-LABEL: @test___drcp_rn(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[DIV_I:%.*]] = fdiv contract double 1.000000e+00, [[X:%.*]]
+// APPROX-NEXT:    ret double [[DIV_I]]
+//
 extern "C" __device__ double test___drcp_rn(double x) {
   return __drcp_rn(x);
 }
@@ -3623,6 +4843,11 @@ extern "C" __device__ double test___drcp_rn(double x) {
 // FINITEONLY-NEXT:    [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_sqrt_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
 // FINITEONLY-NEXT:    ret double [[CALL_I]]
 //
+// APPROX-LABEL: @test___dsqrt_rn(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[CALL_I:%.*]] = tail call contract double @__ocml_sqrt_f64(double noundef [[X:%.*]]) #[[ATTR14]]
+// APPROX-NEXT:    ret double [[CALL_I]]
+//
 extern "C" __device__ double test___dsqrt_rn(double x) {
   return __dsqrt_rn(x);
 }
@@ -3637,6 +4862,11 @@ extern "C" __device__ double test___dsqrt_rn(double x) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
 // FINITEONLY-NEXT:    ret double [[TMP0]]
 //
+// APPROX-LABEL: @test__fma_rn(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
+// APPROX-NEXT:    ret double [[TMP0]]
+//
 extern "C" __device__ double test__fma_rn(double x, double y, double z) {
   return __fma_rn(x, y, z);
 }
@@ -3651,6 +4881,11 @@ extern "C" __device__ double test__fma_rn(double x, double y, double z) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test_float_min(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test_float_min(float x, float y) {
   return min(x, y);
 }
@@ -3665,6 +4900,11 @@ extern "C" __device__ float test_float_min(float x, float y) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
 // FINITEONLY-NEXT:    ret float [[TMP0]]
 //
+// APPROX-LABEL: @test_float_max(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
+// APPROX-NEXT:    ret float [[TMP0]]
+//
 extern "C" __device__ float test_float_max(float x, float y) {
   return max(x, y);
 }
@@ -3679,6 +4919,11 @@ extern "C" __device__ float test_float_max(float x, float y) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
 // FINITEONLY-NEXT:    ret double [[TMP0]]
 //
+// APPROX-LABEL: @test_double_min(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// APPROX-NEXT:    ret double [[TMP0]]
+//
 extern "C" __device__ double test_double_min(double x, double y) {
   return min(x, y);
 }
@@ -3693,6 +4938,11 @@ extern "C" __device__ double test_double_min(double x, double y) {
 // FINITEONLY-NEXT:    [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
 // FINITEONLY-NEXT:    ret double [[TMP0]]
 //
+// APPROX-LABEL: @test_double_max(
+// APPROX-NEXT:  entry:
+// APPROX-NEXT:    [[TMP0:%.*]] = tail call contract double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
+// APPROX-NEXT:    ret double [[TMP0]]
+//
 extern "C" __device__ double test_double_max(double x, double y) {
   return max(x, y);
 }