[libc] Fix missing GPU math implementations (llvm#65616)

These functions were implemented by simply calling their `__builtin_*`
equivalents.
The builtins were resolving to the libc functions back again. This patch
adds explicit
vendor versions for these functions to avoid the recursion.

Author: jhuber6

libc/src/math/gpu/CMakeLists.txt

@@ -183,26 +183,6 @@ add_math_entrypoint_gpu_object(
     -O2
 )
 
-add_math_entrypoint_gpu_object(
-  frexp
-  SRCS
-    frexp.cpp
-  HDRS
-    ../frexp.h
-  COMPILE_OPTIONS
-    -O2
-)
-
-add_math_entrypoint_gpu_object(
-  frexpf
-  SRCS
-    frexpf.cpp
-  HDRS
-    ../frexpf.h
-  COMPILE_OPTIONS
-    -O2
-)
-
 add_math_entrypoint_gpu_object(
   modf
   SRCS
@@ -263,26 +243,6 @@ add_math_entrypoint_gpu_object(
     -O2
 )
 
-add_math_entrypoint_gpu_object(
-  remquo
-  SRCS
-    remquo.cpp
-  HDRS
-    ../remquo.h
-  COMPILE_OPTIONS
-    -O2
-)
-
-add_math_entrypoint_gpu_object(
-  remquof
-  SRCS
-    remquof.cpp
-  HDRS
-    ../remquof.h
-  COMPILE_OPTIONS
-    -O2
-)
-
 add_math_entrypoint_gpu_object(
   rint
   SRCS
@@ -313,26 +273,6 @@ add_math_entrypoint_gpu_object(
     -O2
 )
 
-add_math_entrypoint_gpu_object(
-  scalbn
-  SRCS
-    scalbn.cpp
-  HDRS
-    ../scalbn.h
-  COMPILE_OPTIONS
-    -O2
-)
-
-add_math_entrypoint_gpu_object(
-  scalbnf
-  SRCS
-    scalbnf.cpp
-  HDRS
-    ../scalbnf.h
-  COMPILE_OPTIONS
-    -O2
-)
-
 add_math_entrypoint_gpu_object(
   sinh
   SRCS

libc/src/math/gpu/vendor/CMakeLists.txt

@@ -293,6 +293,29 @@ add_entrypoint_object(
     -O2
 )
 
+add_entrypoint_object(
+  remquo
+  SRCS
+    remquo.cpp
+  HDRS
+    ../../remquo.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+add_entrypoint_object(
+  remquof
+  SRCS
+    remquof.cpp
+  HDRS
+    ../../remquof.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+
 add_entrypoint_object(
   llround
   SRCS
@@ -315,6 +338,29 @@ add_entrypoint_object(
     -O2
 )
 
+add_entrypoint_object(
+  scalbn
+  SRCS
+    scalbn.cpp
+  HDRS
+    ../../scalbn.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+add_entrypoint_object(
+  scalbnf
+  SRCS
+    scalbnf.cpp
+  HDRS
+    ../../scalbnf.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+
 add_entrypoint_object(
   nextafter
   SRCS
@@ -468,3 +514,25 @@ add_entrypoint_object(
     ${bitcode_link_flags}
     -O2
 )
+
+add_entrypoint_object(
+  frexp
+  SRCS
+    frexp.cpp
+  HDRS
+    ../../frexp.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+add_entrypoint_object(
+  frexpf
+  SRCS
+    frexpf.cpp
+  HDRS
+    ../../frexpf.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)

libc/src/math/gpu/vendor/amdgpu/amdgpu.h

@@ -72,6 +72,30 @@ LIBC_INLINE double tan(double x) { return __ocml_tan_f64(x); }
 LIBC_INLINE float tanf(float x) { return __ocml_tan_f32(x); }
 LIBC_INLINE double tanh(double x) { return __ocml_tanh_f64(x); }
 LIBC_INLINE float tanhf(float x) { return __ocml_tanh_f32(x); }
+LIBC_INLINE double scalbn(double x, int i) {
+  return __builtin_amdgcn_ldexp(x, i);
+}
+LIBC_INLINE float scalbnf(float x, int i) {
+  return __builtin_amdgcn_ldexpf(x, i);
+}
+LIBC_INLINE double frexp(double x, int *nptr) {
+  return __builtin_frexp(x, nptr);
+}
+LIBC_INLINE float frexpf(float x, int *nptr) {
+  return __builtin_frexpf(x, nptr);
+}
+LIBC_INLINE double remquo(double x, double y, int *q) {
+  int tmp;
+  double r = __ocml_remquo_f64(x, y, (gpu::Private<int> *)&tmp);
+  *q = tmp;
+  return r;
+}
+LIBC_INLINE float remquof(float x, float y, int *q) {
+  int tmp;
+  float r = __ocml_remquo_f32(x, y, (gpu::Private<int> *)&tmp);
+  *q = tmp;
+  return r;
+}
 
 } // namespace internal
 } // namespace __llvm_libc

libc/src/math/gpu/vendor/amdgpu/declarations.h

@@ -9,6 +9,8 @@
 #ifndef LLVM_LIBC_SRC_MATH_GPU_AMDGPU_DECLARATIONS_H
 #define LLVM_LIBC_SRC_MATH_GPU_AMDGPU_DECLARATIONS_H
 
+#include "src/__support/GPU/utils.h"
+
 namespace __llvm_libc {
 
 extern "C" {
@@ -52,6 +54,8 @@ float __ocml_tan_f32(float);
 double __ocml_tan_f64(double);
 float __ocml_tanh_f32(float);
 double __ocml_tanh_f64(double);
+float __ocml_remquo_f32(float, float, gpu::Private<int> *);
+double __ocml_remquo_f64(double, double, gpu::Private<int> *);
 }
 
 } // namespace __llvm_libc

libc/src/math/gpu/frexp.cpp renamed to libc/src/math/gpu/vendor/frexp.cpp

@@ -9,10 +9,12 @@
 #include "src/math/frexp.h"
 #include "src/__support/common.h"
 
+#include "common.h"
+
 namespace __llvm_libc {
 
 LLVM_LIBC_FUNCTION(double, frexp, (double x, int *p)) {
-  return __builtin_frexp(x, p);
+  return internal::frexp(x, p);
 }
 
 } // namespace __llvm_libc

libc/src/math/gpu/frexpf.cpp renamed to libc/src/math/gpu/vendor/frexpf.cpp

@@ -9,10 +9,12 @@
 #include "src/math/frexpf.h"
 #include "src/__support/common.h"
 
+#include "common.h"
+
 namespace __llvm_libc {
 
 LLVM_LIBC_FUNCTION(float, frexpf, (float x, int *p)) {
-  return __builtin_frexpf(x, p);
+  return internal::frexpf(x, p);
 }
 
 } // namespace __llvm_libc

libc/src/math/gpu/vendor/nvptx/declarations.h

@@ -52,6 +52,12 @@ double __nv_tan(double);
 float __nv_tanf(float);
 double __nv_tanh(double);
 float __nv_tanhf(float);
+double __nv_frexp(double, int *);
+float __nv_frexpf(float, int *);
+double __nv_scalbn(double, int);
+float __nv_scalbnf(float, int);
+double __nv_remquo(double, double, int *);
+float __nv_remquof(float, float, int *);
 }
 
 } // namespace __llvm_libc

libc/src/math/gpu/vendor/nvptx/nvptx.h

@@ -61,6 +61,16 @@ LIBC_INLINE double tan(double x) { return __nv_tan(x); }
 LIBC_INLINE float tanf(float x) { return __nv_tanf(x); }
 LIBC_INLINE double tanh(double x) { return __nv_tanh(x); }
 LIBC_INLINE float tanhf(float x) { return __nv_tanhf(x); }
+LIBC_INLINE double scalbn(double x, int i) { return __nv_scalbn(x, i); }
+LIBC_INLINE float scalbnf(float x, int i) { return __nv_scalbnf(x, i); }
+LIBC_INLINE double frexp(double x, int *i) { return __nv_frexp(x, i); }
+LIBC_INLINE float frexpf(float x, int *i) { return __nv_frexpf(x, i); }
+LIBC_INLINE double remquo(double x, double y, int *i) {
+  return __nv_remquo(x, y, i);
+}
+LIBC_INLINE float remquof(float x, float y, int *i) {
+  return __nv_remquof(x, y, i);
+}
 
 } // namespace internal
 } // namespace __llvm_libc

libc/src/math/gpu/remquo.cpp renamed to libc/src/math/gpu/vendor/remquo.cpp

@@ -9,10 +9,12 @@
 #include "src/math/remquo.h"
 #include "src/__support/common.h"
 
+#include "common.h"
+
 namespace __llvm_libc {
 
 LLVM_LIBC_FUNCTION(double, remquo, (double x, double y, int *quo)) {
-  return __builtin_remquo(x, y, quo);
+  return internal::remquo(x, y, quo);
 }
 
 } // namespace __llvm_libc

libc/src/math/gpu/remquof.cpp renamed to libc/src/math/gpu/vendor/remquof.cpp

@@ -9,10 +9,12 @@
 #include "src/math/remquof.h"
 #include "src/__support/common.h"
 
+#include "common.h"
+
 namespace __llvm_libc {
 
 LLVM_LIBC_FUNCTION(float, remquof, (float x, float y, int *quo)) {
-  return __builtin_remquof(x, y, quo);
+  return internal::remquof(x, y, quo);
 }
 
 } // namespace __llvm_libc

libc/src/math/gpu/scalbn.cpp renamed to libc/src/math/gpu/vendor/scalbn.cpp

@@ -9,10 +9,12 @@
 #include "src/math/scalbn.h"
 #include "src/__support/common.h"
 
+#include "common.h"
+
 namespace __llvm_libc {
 
 LLVM_LIBC_FUNCTION(double, scalbn, (double x, int y)) {
-  return __builtin_scalbn(x, y);
+  return internal::scalbn(x, y);
 }
 
 } // namespace __llvm_libc

libc/src/math/gpu/scalbnf.cpp renamed to libc/src/math/gpu/vendor/scalbnf.cpp

@@ -9,10 +9,12 @@
 #include "src/math/scalbnf.h"
 #include "src/__support/common.h"
 
+#include "common.h"
+
 namespace __llvm_libc {
 
 LLVM_LIBC_FUNCTION(float, scalbnf, (float x, int y)) {
-  return __builtin_scalbnf(x, y);
+  return internal::scalbnf(x, y);
 }
 
 } // namespace __llvm_libc