[libc-commits] [libc] [libc] Fix missing math implementations (PR #65616)
Joseph Huber via libc-commits
libc-commits at lists.llvm.org
Thu Sep 7 07:58:23 PDT 2023
https://github.com/jhuber6 created https://github.com/llvm/llvm-project/pull/65616:
Summary:
These functions were listed as being implemented directly on the GPU,
but they simply reduced to unresolved functions. Move these into the
vendor implementation so they they are provided and do not get used for
tests.
>From ec0e3087c81d7ab86af8b18c86b3c6e054c20b42 Mon Sep 17 00:00:00 2001
From: Joseph Huber <jhuber6 at vols.utk.edu>
Date: Thu, 7 Sep 2023 09:40:10 -0500
Subject: [PATCH] [libc] Fix missing math implementations
Summary:
These functions were listed as being implemented directly on the GPU,
but they simply reduced to unresolved functions. Move these into the
vendor implementation so they they are provided and do not get used for
tests.
---
libc/src/math/gpu/CMakeLists.txt | 60 ----------------
libc/src/math/gpu/vendor/CMakeLists.txt | 68 +++++++++++++++++++
libc/src/math/gpu/vendor/amdgpu/amdgpu.h | 24 +++++++
.../src/math/gpu/vendor/amdgpu/declarations.h | 4 ++
libc/src/math/gpu/{ => vendor}/frexp.cpp | 4 +-
libc/src/math/gpu/{ => vendor}/frexpf.cpp | 4 +-
libc/src/math/gpu/vendor/nvptx/declarations.h | 6 ++
libc/src/math/gpu/vendor/nvptx/nvptx.h | 10 +++
libc/src/math/gpu/{ => vendor}/remquo.cpp | 4 +-
libc/src/math/gpu/{ => vendor}/remquof.cpp | 4 +-
libc/src/math/gpu/{ => vendor}/scalbn.cpp | 4 +-
libc/src/math/gpu/{ => vendor}/scalbnf.cpp | 4 +-
12 files changed, 130 insertions(+), 66 deletions(-)
rename libc/src/math/gpu/{ => vendor}/frexp.cpp (91%)
rename libc/src/math/gpu/{ => vendor}/frexpf.cpp (90%)
rename libc/src/math/gpu/{ => vendor}/remquo.cpp (90%)
rename libc/src/math/gpu/{ => vendor}/remquof.cpp (90%)
rename libc/src/math/gpu/{ => vendor}/scalbn.cpp (90%)
rename libc/src/math/gpu/{ => vendor}/scalbnf.cpp (90%)
diff --git a/libc/src/math/gpu/CMakeLists.txt b/libc/src/math/gpu/CMakeLists.txt
index a23f660261571b..cee7b7d9db476f 100644
--- a/libc/src/math/gpu/CMakeLists.txt
+++ b/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
diff --git a/libc/src/math/gpu/vendor/CMakeLists.txt b/libc/src/math/gpu/vendor/CMakeLists.txt
index a528641a3501b6..2ee74a06a02d46 100644
--- a/libc/src/math/gpu/vendor/CMakeLists.txt
+++ b/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
+)
diff --git a/libc/src/math/gpu/vendor/amdgpu/amdgpu.h b/libc/src/math/gpu/vendor/amdgpu/amdgpu.h
index 3f25dea8b379d0..7755174e445b22 100644
--- a/libc/src/math/gpu/vendor/amdgpu/amdgpu.h
+++ b/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
diff --git a/libc/src/math/gpu/vendor/amdgpu/declarations.h b/libc/src/math/gpu/vendor/amdgpu/declarations.h
index cb910b03b8b261..7219d5a7dfa6d7 100644
--- a/libc/src/math/gpu/vendor/amdgpu/declarations.h
+++ b/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
diff --git a/libc/src/math/gpu/frexp.cpp b/libc/src/math/gpu/vendor/frexp.cpp
similarity index 91%
rename from libc/src/math/gpu/frexp.cpp
rename to libc/src/math/gpu/vendor/frexp.cpp
index eaaf1753a433a8..4bd47fa2810232 100644
--- a/libc/src/math/gpu/frexp.cpp
+++ b/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
diff --git a/libc/src/math/gpu/frexpf.cpp b/libc/src/math/gpu/vendor/frexpf.cpp
similarity index 90%
rename from libc/src/math/gpu/frexpf.cpp
rename to libc/src/math/gpu/vendor/frexpf.cpp
index b600f650da00c2..c6d398007cad4b 100644
--- a/libc/src/math/gpu/frexpf.cpp
+++ b/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
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/nvptx/declarations.h
index 24711d53f75099..8b6702834a04cf 100644
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/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
diff --git a/libc/src/math/gpu/vendor/nvptx/nvptx.h b/libc/src/math/gpu/vendor/nvptx/nvptx.h
index 3310d2406c9d54..6ea1743cf7a6f3 100644
--- a/libc/src/math/gpu/vendor/nvptx/nvptx.h
+++ b/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
diff --git a/libc/src/math/gpu/remquo.cpp b/libc/src/math/gpu/vendor/remquo.cpp
similarity index 90%
rename from libc/src/math/gpu/remquo.cpp
rename to libc/src/math/gpu/vendor/remquo.cpp
index e8f55231950ad3..755efeaba530ba 100644
--- a/libc/src/math/gpu/remquo.cpp
+++ b/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
diff --git a/libc/src/math/gpu/remquof.cpp b/libc/src/math/gpu/vendor/remquof.cpp
similarity index 90%
rename from libc/src/math/gpu/remquof.cpp
rename to libc/src/math/gpu/vendor/remquof.cpp
index 310734756abf83..c2be2cd87aad6c 100644
--- a/libc/src/math/gpu/remquof.cpp
+++ b/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
diff --git a/libc/src/math/gpu/scalbn.cpp b/libc/src/math/gpu/vendor/scalbn.cpp
similarity index 90%
rename from libc/src/math/gpu/scalbn.cpp
rename to libc/src/math/gpu/vendor/scalbn.cpp
index 99f67a9bfe513b..44aa286166cbf5 100644
--- a/libc/src/math/gpu/scalbn.cpp
+++ b/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
diff --git a/libc/src/math/gpu/scalbnf.cpp b/libc/src/math/gpu/vendor/scalbnf.cpp
similarity index 90%
rename from libc/src/math/gpu/scalbnf.cpp
rename to libc/src/math/gpu/vendor/scalbnf.cpp
index 6e6915b51bbda9..6d02f8617201ac 100644
--- a/libc/src/math/gpu/scalbnf.cpp
+++ b/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
More information about the libc-commits
mailing list