diff --git a/libc/config/gpu/entrypoints.txt b/libc/config/gpu/entrypoints.txt
--- a/libc/config/gpu/entrypoints.txt
+++ b/libc/config/gpu/entrypoints.txt
@@ -83,8 +83,55 @@
 
 set(TARGET_LIBM_ENTRYPOINTS
     # math.h entrypoints
+    libc.src.math.acosf
+    libc.src.math.acoshf
+    libc.src.math.asinf
+    libc.src.math.asinhf
+    libc.src.math.atanf
+    libc.src.math.atanhf
+    libc.src.math.ceil
+    libc.src.math.ceilf
+    libc.src.math.ceill
+    libc.src.math.copysign
+    libc.src.math.copysignf
+    libc.src.math.copysignl
+    libc.src.math.cos
+    libc.src.math.cosf
+    libc.src.math.coshf
+    libc.src.math.exp10f
+    libc.src.math.exp2f
+    libc.src.math.expf
+    libc.src.math.expm1f
+    libc.src.math.fabs
+    libc.src.math.fabsf
+    libc.src.math.fabsl
+    libc.src.math.fdim
+    libc.src.math.fdimf
+    libc.src.math.floor
+    libc.src.math.floorf
+    libc.src.math.floorl
     libc.src.math.fma
     libc.src.math.fmaf
+    libc.src.math.fmax
+    libc.src.math.fmaxf
+    libc.src.math.fmaxl
+    libc.src.math.fmin
+    libc.src.math.fminf
+    libc.src.math.fminl
+    libc.src.math.fmod
+    libc.src.math.fmodf
+    libc.src.math.frexp
+    libc.src.math.frexpf
+    libc.src.math.hypot
+    libc.src.math.hypotf
+    libc.src.math.ilogb
+    libc.src.math.ilogbf
+    libc.src.math.ldexp
+    libc.src.math.ldexpf
+    libc.src.math.llrint
+    libc.src.math.llrintf
+    libc.src.math.llround
+    libc.src.math.llroundf
     libc.src.math.sin
     libc.src.math.round
     libc.src.math.roundf
diff --git a/libc/src/math/gpu/CMakeLists.txt b/libc/src/math/gpu/CMakeLists.txt
--- a/libc/src/math/gpu/CMakeLists.txt
+++ b/libc/src/math/gpu/CMakeLists.txt
@@ -23,6 +23,306 @@
   )
 endfunction()
 
+add_math_entrypoint_gpu_object(
+  ceil
+  SRCS
+    ceil.cpp
+  HDRS
+    ../ceil.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  ceilf
+  SRCS
+    ceilf.cpp
+  HDRS
+    ../ceilf.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  ceill
+  SRCS
+    ceill.cpp
+  HDRS
+    ../ceill.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  copysign
+  SRCS
+    copysign.cpp
+  HDRS
+    ../copysign.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  copysignf
+  SRCS
+    copysignf.cpp
+  HDRS
+    ../copysignf.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  copysignl
+  SRCS
+    copysignl.cpp
+  HDRS
+    ../copysignl.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  cos
+  SRCS
+    cos.cpp
+  HDRS
+    ../cos.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  cosf
+  SRCS
+    cosf.cpp
+  HDRS
+    ../cosf.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  exp2f
+  SRCS
+    exp2f.cpp
+  HDRS
+    ../exp2f.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  expf
+  SRCS
+    expf.cpp
+  HDRS
+    ../expf.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  fabs
+  SRCS
+    fabs.cpp
+  HDRS
+    ../fabs.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  fabsf
+  SRCS
+    fabsf.cpp
+  HDRS
+    ../fabsf.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  fabsl
+  SRCS
+    fabsl.cpp
+  HDRS
+    ../fabsl.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  floor
+  SRCS
+    floor.cpp
+  HDRS
+    ../floor.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  floorf
+  SRCS
+    floorf.cpp
+  HDRS
+    ../floorf.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  floorl
+  SRCS
+    floorl.cpp
+  HDRS
+    ../floorl.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  fma
+  SRCS
+    fma.cpp
+  HDRS
+    ../fma.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  fmaf
+  SRCS
+    fmaf.cpp
+  HDRS
+    ../fmaf.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  fmax
+  SRCS
+    fmax.cpp
+  HDRS
+    ../fmax.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  fmaxf
+  SRCS
+    fmaxf.cpp
+  HDRS
+    ../fmaxf.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  fmaxl
+  SRCS
+    fmaxl.cpp
+  HDRS
+    ../fmaxl.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  fmin
+  SRCS
+    fmin.cpp
+  HDRS
+    ../fmin.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  fminf
+  SRCS
+    fminf.cpp
+  HDRS
+    ../fminf.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  fminl
+  SRCS
+    fminl.cpp
+  HDRS
+    ../fminl.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  fmod
+  SRCS
+    fmod.cpp
+  HDRS
+    ../fmod.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  fmodf
+  SRCS
+    fmodf.cpp
+  HDRS
+    ../fmodf.h
+  COMPILE_OPTIONS
+    -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(
+  ldexp
+  SRCS
+    ldexp.cpp
+  HDRS
+    ../ldexp.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_math_entrypoint_gpu_object(
+  ldexpf
+  SRCS
+    ldexpf.cpp
+  HDRS
+    ../ldexpf.h
+  COMPILE_OPTIONS
+    -O2
+)
+
 add_math_entrypoint_gpu_object(
   round
   SRCS
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/ceil.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/ceil.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/ceil.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the ceil function for GPU -----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/ceil.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(double, ceil, (double x)) { return __builtin_ceil(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/ceilf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/ceilf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/ceilf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the ceilf function for GPU ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/ceilf.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(float, ceilf, (float x)) { return __builtin_ceilf(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/amdgpu/declarations.h b/libc/src/math/gpu/ceill.cpp
copy from libc/src/math/gpu/vendor/amdgpu/declarations.h
copy to libc/src/math/gpu/ceill.cpp
--- a/libc/src/math/gpu/vendor/amdgpu/declarations.h
+++ b/libc/src/math/gpu/ceill.cpp
@@ -1,4 +1,4 @@
-//===-- AMDGPU specific declarations for math support ---------------------===//
+//===-- Implementation of the ceill function for GPU ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,18 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_AMDGPU_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_AMDGPU_DECLARATIONS_H
+#include "src/math/ceill.h"
+#include "src/__support/FPUtil/PlatformDefs.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __ocml_sin_f64(double);
+#ifndef LONG_DOUBLE_IS_DOUBLE
+#error "GPU targets do not support long doubles"
+#endif
+
+LLVM_LIBC_FUNCTION(long double, ceill, (long double x)) {
+  return __builtin_ceill(x);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_AMDGPU_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/copysign.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/copysign.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/copysign.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the copysign function for GPU -------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/copysign.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(double, copysign, (double x, double y)) {
+  return __builtin_copysign(x, y);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/copysignf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/copysignf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/copysignf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the copysignf function for GPU ------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/copysignf.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(float, copysignf, (float x, float y)) {
+  return __builtin_copysignf(x, y);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/copysignl.cpp b/libc/src/math/gpu/copysignl.cpp
new file mode 100644
--- /dev/null
+++ b/libc/src/math/gpu/copysignl.cpp
@@ -0,0 +1,23 @@
+//===-- Implementation of the copysignl function for GPU ------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/copysignl.h"
+#include "src/__support/FPUtil/PlatformDefs.h"
+#include "src/__support/common.h"
+
+namespace __llvm_libc {
+
+#ifndef LONG_DOUBLE_IS_DOUBLE
+#error "GPU targets do not support long doubles"
+#endif
+
+LLVM_LIBC_FUNCTION(long double, copysignl, (long double x, long double y)) {
+  return __builtin_copysignl(x, y);
+}
+
+} // namespace __llvm_libc
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/cos.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/cos.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/cos.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the cos function for GPU ------------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/cos.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(double, cos, (double x)) { return __builtin_cos(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/cosf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/cosf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/cosf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the cosf function for GPU -----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/cosf.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(float, cosf, (float x)) { return __builtin_cosf(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/exp2f.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/exp2f.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/exp2f.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the exp2f function for GPU ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/exp2f.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(float, exp2f, (float x)) { return __builtin_exp2f(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/expf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/expf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/expf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the expf function for GPU -----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/expf.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(float, expf, (float x)) { return __builtin_expf(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/fabs.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/fabs.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/fabs.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the fabs function for GPU -----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/fabs.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(double, fabs, (double x)) { return __builtin_fabs(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/fabsf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/fabsf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/fabsf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the fabsf function for GPU ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/fabsf.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(float, fabsf, (float x)) { return __builtin_fabsf(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/amdgpu/declarations.h b/libc/src/math/gpu/fabsl.cpp
copy from libc/src/math/gpu/vendor/amdgpu/declarations.h
copy to libc/src/math/gpu/fabsl.cpp
--- a/libc/src/math/gpu/vendor/amdgpu/declarations.h
+++ b/libc/src/math/gpu/fabsl.cpp
@@ -1,4 +1,4 @@
-//===-- AMDGPU specific declarations for math support ---------------------===//
+//===-- Implementation of the fabsl function for GPU ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,18 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_AMDGPU_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_AMDGPU_DECLARATIONS_H
+#include "src/math/fabsl.h"
+#include "src/__support/FPUtil/PlatformDefs.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __ocml_sin_f64(double);
+#ifndef LONG_DOUBLE_IS_DOUBLE
+#error "GPU targets do not support long doubles"
+#endif
+
+LLVM_LIBC_FUNCTION(long double, fabsl, (long double x)) {
+  return __builtin_fabsl(x);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_AMDGPU_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/floor.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/floor.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/floor.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the floor function for GPU ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/floor.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(double, floor, (double x)) { return __builtin_floor(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/floorf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/floorf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/floorf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the floorf function for GPU ---------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,11 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/floorf.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(float, floorf, (float x)) { return __builtin_floorf(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/amdgpu/declarations.h b/libc/src/math/gpu/floorl.cpp
copy from libc/src/math/gpu/vendor/amdgpu/declarations.h
copy to libc/src/math/gpu/floorl.cpp
--- a/libc/src/math/gpu/vendor/amdgpu/declarations.h
+++ b/libc/src/math/gpu/floorl.cpp
@@ -1,4 +1,4 @@
-//===-- AMDGPU specific declarations for math support ---------------------===//
+//===-- Implementation of the floorl function for GPU ---------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,18 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_AMDGPU_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_AMDGPU_DECLARATIONS_H
+#include "src/math/floorl.h"
+#include "src/__support/FPUtil/PlatformDefs.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __ocml_sin_f64(double);
+#ifndef LONG_DOUBLE_IS_DOUBLE
+#error "GPU targets do not support long doubles"
+#endif
+
+LLVM_LIBC_FUNCTION(long double, floorl, (long double x)) {
+  return __builtin_floorl(x);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_AMDGPU_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/fma.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/fma.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/fma.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the fma function for GPU ------------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/fma.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(double, fma, (double x, double y, double z)) {
+  return __builtin_fma(x, y, z);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/fmaf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/fmaf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/fmaf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the fmaf function for GPU -----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/fmaf.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(float, fmaf, (float x, float y, float z)) {
+  return __builtin_fmaf(x, y, z);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/fmax.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/fmax.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/fmax.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the fmax function for GPU -----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/fmax.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(double, fmax, (double x, double y)) {
+  return __builtin_fmax(x, y);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/fmaxf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/fmaxf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/fmaxf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the fmaxf function for GPU ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/fmaxf.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(float, fmaxf, (float x, float y)) {
+  return __builtin_fmaxf(x, y);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/fmaxl.cpp b/libc/src/math/gpu/fmaxl.cpp
new file mode 100644
--- /dev/null
+++ b/libc/src/math/gpu/fmaxl.cpp
@@ -0,0 +1,23 @@
+//===-- Implementation of the fmaxl function for GPU ----------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/fmaxl.h"
+#include "src/__support/FPUtil/PlatformDefs.h"
+#include "src/__support/common.h"
+
+namespace __llvm_libc {
+
+#ifndef LONG_DOUBLE_IS_DOUBLE
+#error "GPU targets do not support long doubles"
+#endif
+
+LLVM_LIBC_FUNCTION(long double, fmaxl, (long double x, long double y)) {
+  return __builtin_fmaxl(x, y);
+}
+
+} // namespace __llvm_libc
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/fmin.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/fmin.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/fmin.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the fmin function for GPU -----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/fmin.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(double, fmin, (double x, double y)) {
+  return __builtin_fmin(x, y);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/fminf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/fminf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/fminf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the fminf function for GPU ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/fminf.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(float, fminf, (float x, float y)) {
+  return __builtin_fminf(x, y);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/fminl.cpp b/libc/src/math/gpu/fminl.cpp
new file mode 100644
--- /dev/null
+++ b/libc/src/math/gpu/fminl.cpp
@@ -0,0 +1,23 @@
+//===-- Implementation of the fminl function for GPU ----------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/fminl.h"
+#include "src/__support/FPUtil/PlatformDefs.h"
+#include "src/__support/common.h"
+
+namespace __llvm_libc {
+
+#ifndef LONG_DOUBLE_IS_DOUBLE
+#error "GPU targets do not support long doubles"
+#endif
+
+LLVM_LIBC_FUNCTION(long double, fminl, (long double x, long double y)) {
+  return __builtin_fminl(x, y);
+}
+
+} // namespace __llvm_libc
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/fmod.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/fmod.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/fmod.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the fmod function for GPU -----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/fmod.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(double, fmod, (double x, double y)) {
+  return __builtin_fmod(x, y);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/fmodf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/fmodf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/fmodf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the fmodf function for GPU ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/fmodf.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(float, fmodf, (float x, float y)) {
+  return __builtin_fmodf(x, y);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/frexp.cpp b/libc/src/math/gpu/frexp.cpp
new file mode 100644
--- /dev/null
+++ b/libc/src/math/gpu/frexp.cpp
@@ -0,0 +1,30 @@
+//===-- Implementation of the frexp function for GPU ----------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/frexp.h"
+#include "src/__support/common.h"
+#include "src/__support/macros/properties/architectures.h"
+
+#if defined(LIBC_TARGET_ARCH_IS_NVPTX)
+#include "nvptx/nvptx.h"
+#endif
+
+namespace __llvm_libc {
+
+LLVM_LIBC_FUNCTION(double, frexp, (double x, int *p)) {
+#if defined(LIBC_TARGET_ARCH_IS_NVPTX)
+  return __nv_frexp(x, p);
+#elif defined(LIBC_TARGET_ARCH_IS_AMDGPU)
+  *p = __builtin_amdgcn_frexp_exp(x);
+  return __builtin_amdgcn_frexp_mant(x);
+#else
+#error "Unsupported platform"
+#endif
+}
+
+} // namespace __llvm_libc
diff --git a/libc/src/math/gpu/frexpf.cpp b/libc/src/math/gpu/frexpf.cpp
new file mode 100644
--- /dev/null
+++ b/libc/src/math/gpu/frexpf.cpp
@@ -0,0 +1,30 @@
+//===-- Implementation of the frexpf function for GPU ---------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/frexpf.h"
+#include "src/__support/common.h"
+#include "src/__support/macros/properties/architectures.h"
+
+#if defined(LIBC_TARGET_ARCH_IS_NVPTX)
+#include "nvptx/nvptx.h"
+#endif
+
+namespace __llvm_libc {
+
+LLVM_LIBC_FUNCTION(float, frexpf, (float x, int *p)) {
+#if defined(LIBC_TARGET_ARCH_IS_NVPTX)
+  return __nv_frexpf(x, p);
+#elif defined(LIBC_TARGET_ARCH_IS_AMDGPU)
+  *p = __builtin_amdgcn_frexp_expf(x);
+  return __builtin_amdgcn_frexp_mantf(x);
+#else
+#error "Unsupported platform"
+#endif
+}
+
+} // namespace __llvm_libc
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/ldexp.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/ldexp.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/ldexp.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the ldexp function for GPU ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/ldexp.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(double, ldexp, (double x, int y)) {
+  return __builtin_ldexp(x, y);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/ldexpf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/ldexpf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/ldexpf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the ldexpf function for GPU ---------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/ldexpf.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(float, ldexpf, (float x, int y)) {
+  return __builtin_ldexpf(x, y);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/CMakeLists.txt b/libc/src/math/gpu/vendor/CMakeLists.txt
--- a/libc/src/math/gpu/vendor/CMakeLists.txt
+++ b/libc/src/math/gpu/vendor/CMakeLists.txt
@@ -29,6 +29,212 @@
 # will link in identity metadata from both libraries. This silences the warning.
 list(APPEND bitcode_link_flags "-Wno-linker-warnings")
 
+add_entrypoint_object(
+  acosf
+  SRCS
+    acosf.cpp
+  HDRS
+    ../../acosf.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+add_entrypoint_object(
+  acoshf
+  SRCS
+    acoshf.cpp
+  HDRS
+    ../../acoshf.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+add_entrypoint_object(
+  asinf
+  SRCS
+    asinf.cpp
+  HDRS
+    ../../asinf.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+add_entrypoint_object(
+  asinhf
+  SRCS
+    asinhf.cpp
+  HDRS
+    ../../asinhf.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+add_entrypoint_object(
+  atanf
+  SRCS
+    atanf.cpp
+  HDRS
+    ../../atanf.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+add_entrypoint_object(
+  atanhf
+  SRCS
+    atanhf.cpp
+  HDRS
+    ../../atanhf.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+add_entrypoint_object(
+  coshf
+  SRCS
+    coshf.cpp
+  HDRS
+    ../../coshf.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+add_entrypoint_object(
+  exp10f
+  SRCS
+    exp10f.cpp
+  HDRS
+    ../../exp10f.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+add_entrypoint_object(
+  expm1f
+  SRCS
+    expm1f.cpp
+  HDRS
+    ../../expm1f.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+
+add_entrypoint_object(
+  fdim
+  SRCS
+    fdim.cpp
+  HDRS
+    ../../fdim.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+add_entrypoint_object(
+  fdimf
+  SRCS
+    fdimf.cpp
+  HDRS
+    ../../fdimf.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+add_entrypoint_object(
+  hypot
+  SRCS
+    hypot.cpp
+  HDRS
+    ../../hypot.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+add_entrypoint_object(
+  hypotf
+  SRCS
+    hypotf.cpp
+  HDRS
+    ../../hypotf.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+add_entrypoint_object(
+  ilogb
+  SRCS
+    ilogb.cpp
+  HDRS
+    ../../ilogb.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+add_entrypoint_object(
+  ilogbf
+  SRCS
+    ilogbf.cpp
+  HDRS
+    ../../ilogbf.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+add_entrypoint_object(
+  llrint
+  SRCS
+    llrint.cpp
+  HDRS
+    ../../llrint.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_entrypoint_object(
+  llrintf
+  SRCS
+    llrintf.cpp
+  HDRS
+    ../../llrintf.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_entrypoint_object(
+  llround
+  SRCS
+    llround.cpp
+  HDRS
+    ../../llround.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_entrypoint_object(
+  llroundf
+  SRCS
+    llroundf.cpp
+  HDRS
+    ../../llroundf.h
+  COMPILE_OPTIONS
+    -O2
+)
+
 add_entrypoint_object(
   sin
   SRCS
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/acosf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/acosf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/acosf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the acosf function for GPU ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/acosf.h"
+#include "src/__support/common.h"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(float, acosf, (float x)) { return internal::acosf(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/acoshf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/acoshf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/acoshf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the acoshf function for GPU ---------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/acoshf.h"
+#include "src/__support/common.h"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(float, acoshf, (float x)) { return internal::acoshf(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/amdgpu/amdgpu.h b/libc/src/math/gpu/vendor/amdgpu/amdgpu.h
--- a/libc/src/math/gpu/vendor/amdgpu/amdgpu.h
+++ b/libc/src/math/gpu/vendor/amdgpu/amdgpu.h
@@ -16,7 +16,25 @@
 
 namespace __llvm_libc {
 namespace internal {
-
+LIBC_INLINE float acosf(float x) { return __ocml_acos_f32(x); }
+LIBC_INLINE float acoshf(float x) { return __ocml_acosh_f32(x); }
+LIBC_INLINE float asinf(float x) { return __ocml_asin_f32(x); }
+LIBC_INLINE float asinhf(float x) { return __ocml_asinh_f32(x); }
+LIBC_INLINE float atanf(float x) { return __ocml_atan_f32(x); }
+LIBC_INLINE float atanhf(float x) { return __ocml_atanh_f32(x); }
+LIBC_INLINE float coshf(float x) { return __ocml_cosh_f32(x); }
+LIBC_INLINE float exp10f(float x) { return __ocml_exp10_f32(x); }
+LIBC_INLINE float expm1f(float x) { return __ocml_expm1_f32(x); }
+LIBC_INLINE double fdim(double x, double y) { return __ocml_fdim_f64(x, y); }
+LIBC_INLINE float fdimf(float x, float y) { return __ocml_fdim_f32(x, y); }
+LIBC_INLINE double hypot(double x, double y) { return __ocml_hypot_f64(x, y); }
+LIBC_INLINE float hypotf(float x, float y) { return __ocml_hypot_f32(x, y); }
+LIBC_INLINE int ilogb(double x) { return __ocml_ilogb_f64(x); }
+LIBC_INLINE int ilogbf(float x) { return __ocml_ilogb_f32(x); }
+LIBC_INLINE long long llrint(double x) { return __ocml_rint_f64(x); }
+LIBC_INLINE long long llrintf(float x) { return __ocml_rint_f32(x); }
+LIBC_INLINE long long llround(double x) { return __ocml_round_f64(x); }
+LIBC_INLINE long long llroundf(float x) { return __ocml_round_f32(x); }
 LIBC_INLINE double sin(double x) { return __ocml_sin_f64(x); }
 
 } // namespace internal
diff --git a/libc/src/math/gpu/vendor/amdgpu/declarations.h b/libc/src/math/gpu/vendor/amdgpu/declarations.h
--- a/libc/src/math/gpu/vendor/amdgpu/declarations.h
+++ b/libc/src/math/gpu/vendor/amdgpu/declarations.h
@@ -12,6 +12,25 @@
 namespace __llvm_libc {
 
 extern "C" {
+float __ocml_acos_f32(float);
+float __ocml_acosh_f32(float);
+float __ocml_asin_f32(float);
+float __ocml_asinh_f32(float);
+float __ocml_atan_f32(float);
+float __ocml_atanh_f32(float);
+float __ocml_cosh_f32(float);
+float __ocml_exp10_f32(float);
+float __ocml_expm1_f32(float);
+float __ocml_fdim_f32(float, float);
+double __ocml_fdim_f64(double, double);
+double __ocml_hypot_f64(double, double);
+float __ocml_hypot_f32(float, float);
+int __ocml_ilogb_f64(double);
+int __ocml_ilogb_f32(float);
+double __ocml_rint_f64(double);
+float __ocml_rint_f32(float);
+double __ocml_round_f64(double);
+float __ocml_round_f32(float);
 double __ocml_sin_f64(double);
 }
 
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/asinf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/asinf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/asinf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the asinf function for GPU ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/asinf.h"
+#include "src/__support/common.h"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(float, asinf, (float x)) { return internal::asinf(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/asinhf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/asinhf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/asinhf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the asinhf function for GPU ---------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/asinhf.h"
+#include "src/__support/common.h"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(float, asinhf, (float x)) { return internal::asinhf(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/atanf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/atanf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/atanf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the atanf function for GPU ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/atanf.h"
+#include "src/__support/common.h"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(float, atanf, (float x)) { return internal::atanf(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/atanhf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/atanhf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/atanhf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the atanhf function for GPU ---------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/atanhf.h"
+#include "src/__support/common.h"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(float, atanhf, (float x)) { return internal::atanhf(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/coshf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/coshf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/coshf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the coshf function for GPU ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/coshf.h"
+#include "src/__support/common.h"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(float, coshf, (float x)) { return internal::coshf(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/exp10f.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/exp10f.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/exp10f.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the exp10f function for GPU ---------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/exp10f.h"
+#include "src/__support/common.h"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(float, exp10f, (float x)) { return internal::exp10f(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/expm1f.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/expm1f.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/expm1f.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the expm1f function for GPU ---------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/expm1f.h"
+#include "src/__support/common.h"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(float, expm1f, (float x)) { return internal::expm1f(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/fdim.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/fdim.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/fdim.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the fdim function for GPU -----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/fdim.h"
+#include "src/__support/common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(double, fdim, (double x, double y)) {
+  return __builtin_fdim(x, y);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/fdimf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/fdimf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/fdimf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the fdimf function for GPU ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,15 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/fdimf.h"
+#include "src/__support/common.h"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(float, fdimf, (float x, float y)) {
+  return internal::fdimf(x, y);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/hypot.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/hypot.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/hypot.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the hypot function for GPU ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,15 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/hypot.h"
+#include "src/__support/common.h"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(double, hypot, (double x, double y)) {
+  return internal::hypot(x, y);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/hypotf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/hypotf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/hypotf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the hypotf function for GPU ---------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,15 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/hypotf.h"
+#include "src/__support/common.h"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(float, hypotf, (float x, float y)) {
+  return internal::hypot(x, y);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/ilogb.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/ilogb.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/ilogb.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the ilogb function for GPU ----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/ilogb.h"
+#include "src/__support/common.h"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(int, ilogb, (double x)) { return internal::ilogb(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/ilogbf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/ilogbf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/ilogbf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the ilogbf function for GPU ---------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,13 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/ilogbf.h"
+#include "src/__support/common.h"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(int, ilogbf, (float x)) { return internal::ilogbf(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/llrint.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/llrint.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/llrint.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the llrint function for GPU ---------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,15 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/llrint.h"
+#include "src/__support/common.h"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(long long, llrint, (double x)) {
+  return internal::llrint(x);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/llrintf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/llrintf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/llrintf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the llrintf function for GPU --------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,15 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/llrintf.h"
+#include "src/__support/common.h"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(long long, llrintf, (float x)) {
+  return internal::llrintf(x);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/llround.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/llround.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/llround.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the llround function for GPU --------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,15 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/llround.h"
+#include "src/__support/common.h"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(long long, llround, (double x)) {
+  return internal::llround(x);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/llroundf.cpp
copy from libc/src/math/gpu/vendor/nvptx/declarations.h
copy to libc/src/math/gpu/vendor/llroundf.cpp
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/llroundf.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTX specific declarations for math support ----------------------===//
+//===-- Implementation of the llroundf function for GPU -------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,15 +6,15 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
-#define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
+#include "src/math/llroundf.h"
+#include "src/__support/common.h"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(long long, llroundf, (float x)) {
+  return internal::llroundf(x);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/gpu/vendor/nvptx/declarations.h
--- a/libc/src/math/gpu/vendor/nvptx/declarations.h
+++ b/libc/src/math/gpu/vendor/nvptx/declarations.h
@@ -12,6 +12,27 @@
 namespace __llvm_libc {
 
 extern "C" {
+float __nv_acosf(float);
+float __nv_acoshf(float);
+float __nv_asinf(float);
+float __nv_asinhf(float);
+float __nv_atanf(float);
+float __nv_atanhf(float);
+float __nv_coshf(float);
+float __nv_exp10f(float);
+float __nv_expm1f(float);
+double __nv_fdim(double, double);
+float __nv_fdimf(float, float);
+double __nv_frexp(double, int *);
+float __nv_frexpf(float, int *);
+double __nv_hypot(double, double);
+float __nv_hypotf(float, float);
+int __nv_ilogb(double);
+int __nv_ilogbf(float);
+long long __nv_llrint(double);
+long long __nv_llrintf(float);
+long long __nv_llround(double);
+long long __nv_llroundf(float);
 double __nv_sin(double);
 }
 
diff --git a/libc/src/math/gpu/vendor/nvptx/nvptx.h b/libc/src/math/gpu/vendor/nvptx/nvptx.h
--- a/libc/src/math/gpu/vendor/nvptx/nvptx.h
+++ b/libc/src/math/gpu/vendor/nvptx/nvptx.h
@@ -15,7 +15,25 @@
 
 namespace __llvm_libc {
 namespace internal {
-
+LIBC_INLINE float acosf(float x) { return __nv_acosf(x); }
+LIBC_INLINE float acoshf(float x) { return __nv_acoshf(x); }
+LIBC_INLINE float asinf(float x) { return __nv_asinf(x); }
+LIBC_INLINE float asinhf(float x) { return __nv_asinhf(x); }
+LIBC_INLINE float atanf(float x) { return __nv_atanf(x); }
+LIBC_INLINE float atanhf(float x) { return __nv_atanhf(x); }
+LIBC_INLINE float coshf(float x) { return __nv_coshf(x); }
+LIBC_INLINE float exp10f(float x) { return __nv_exp10f(x); }
+LIBC_INLINE float expm1f(float x) { return __nv_expm1f(x); }
+LIBC_INLINE double fdim(double x, double y) { return __nv_fdim(x, y); }
+LIBC_INLINE float fdimf(float x, float y) { return __nv_fdimf(x, y); }
+LIBC_INLINE double hypot(double x, double y) { return __nv_hypot(x, y); }
+LIBC_INLINE float hypotf(float x, float y) { return __nv_hypotf(x, y); }
+LIBC_INLINE int ilogb(double x) { return __nv_ilogb(x); }
+LIBC_INLINE int ilogbf(float x) { return __nv_ilogbf(x); }
+LIBC_INLINE long long llrint(double x) { return __nv_llrint(x); }
+LIBC_INLINE long long llrintf(float x) { return __nv_llrintf(x); }
+LIBC_INLINE long long llround(double x) { return __nv_llround(x); }
+LIBC_INLINE long long llroundf(float x) { return __nv_llroundf(x); }
 LIBC_INLINE double sin(double x) { return __nv_sin(x); }
 
 } // namespace internal