clang/HIP: Inline frexp/frexpf implementations

Don't bother calling ocml. This stops setting the
appropriate fast math flags, and requires this junk
for passing to a private pointer.

Author: arsenm

clang/lib/Headers/__clang_hip_math.h

@@ -257,15 +257,8 @@ float fmodf(float __x, float __y) { return __ocml_fmod_f32(__x, __y); }
 
 __DEVICE__
 float frexpf(float __x, int *__nptr) {
-  int __tmp;
-#ifdef __OPENMP_AMDGCN__
-#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
-#endif
-  float __r =
-      __ocml_frexp_f32(__x, (__attribute__((address_space(5))) int *)&__tmp);
-  *__nptr = __tmp;
-
-  return __r;
+  *__nptr = __builtin_amdgcn_frexp_expf(__x);
+  return __builtin_amdgcn_frexp_mantf(__x);
 }
 
 __DEVICE__
@@ -813,14 +806,8 @@ double fmod(double __x, double __y) { return __ocml_fmod_f64(__x, __y); }
 
 __DEVICE__
 double frexp(double __x, int *__nptr) {
-  int __tmp;
-#ifdef __OPENMP_AMDGCN__
-#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
-#endif
-  double __r =
-      __ocml_frexp_f64(__x, (__attribute__((address_space(5))) int *)&__tmp);
-  *__nptr = __tmp;
-  return __r;
+  *__nptr = __builtin_amdgcn_frexp_exp(__x);
+  return __builtin_amdgcn_frexp_mant(__x);
 }
 
 __DEVICE__