[libc-commits] [libc] [libc] Fix missing GPU math implementations (PR #65616)
Joseph Huber via libc-commits
libc-commits at lists.llvm.org
Thu Sep 7 09:32:51 PDT 2023
https://github.com/jhuber6 updated https://github.com/llvm/llvm-project/pull/65616:
>From c497acaa88758d0a8a937bfee2ab05de297e4c52 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 GPU vendor math implementations
Summary:
This patch addresses a few math functions that were mistakenly
implemented using only a `__builtin_` variant. These functions appeared
to have been implemented, but were simply lowered to an unresolved
function call by the backend. E.g. `__builtin_xyz` was becoming `xyz`
and would fail to link. This patch moves these functions to the `vendor`
directory where they can me implemented using the external source code.
This also allows math tests to run.
---
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 a23f660261571b2..cee7b7d9db476f2 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 a528641a3501b69..2ee74a06a02d461 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 3f25dea8b379d00..7755174e445b222 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 cb910b03b8b2610..7219d5a7dfa6d73 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 eaaf1753a433a80..4bd47fa2810232a 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 b600f650da00c2b..c6d398007cad4b7 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 24711d53f750998..8b6702834a04cf4 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 3310d2406c9d54f..6ea1743cf7a6f3e 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 e8f55231950ad3c..755efeaba530ba0 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 310734756abf83c..c2be2cd87aad6c6 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 99f67a9bfe513b4..44aa286166cbf58 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 6e6915b51bbda9b..6d02f8617201ac6 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