[libc] Add math functions to AMD/NVPTX libm

[elmcdonough]

Summary

Related to D152486. The following functions are included in this revision: acosf, acoshf, asinf, asinhf, atanf, atanhf, ceil, ceilf, copysign, copysignf, cos, cosf, cosh, coshf, exp10f, exp2f, expf, expm1f, fabs, fabsf, fdim, fdimf, floor, floorf, fma, fmaf, fmax, fmaxf, fmin, fminf, fmod, fmodf, frexp, frexpf, hypot, hypotf, ilogb, ilogbf, ldexp, ldexpf, llrint, llrintf, llround, llroundf, pow, and powf.

[AntonRydahl]

This comment was removed by AntonRydahl.

[elmcdonough]

I have been able to confirm that this patch builds and that each builtin function collapses to an LLVM intrinsic.

That said, when I compile the individual cpp files with llc, I do get an issue related to cos in the 'AMDGPU DAG->DAG Pattern Instruction Selection' pass. Not sure if this is specific to my setup or if cos needs to be a vendor function.

[sivachandra]

libc/src/math/gpu/cos.cpp

14 ↗

(On Diff #533188)

I am not sure the GPUs would have instructions to do trignometric functions.

libc/src/math/gpu/exp2f.cpp

14 ↗

(On Diff #533188)

AMD gpus seem to be having exp2f. Can we assume other GPUs also have them? Even if they do have, what guarantees and promises they come with? The reason I am asking is to understand how we should set up tests for them.

[jhuber6]

libc/src/math/gpu/cos.cpp

14 ↗

(On Diff #533188)

They do not, https://godbolt.org/z/zTWY168a3.

libc/src/math/gpu/exp2f.cpp

14 ↗

(On Diff #533188)

Seems NVPTX does not support it https://godbolt.org/z/a1jGzeo48. Maybe @arsenm can answer the later questions about the properties.

I'm still not sure how to stand up tests for this math on the GPU. We may need to map it against some host implementation and compare within some ULP. We could potentially have the GPU allocate a table like arr = (x, f(x)) for that math function and pass it in, letting the GPU scan it.

[arsenm]

libc/src/math/gpu/cos.cpp

14 ↗

(On Diff #533188)

They do for f32. llvm.sin.f32 and llvm.cos.f32 expand fast-mathy for AMDGPU and I've been debating dropping support for them because it's just wrong

libc/src/math/gpu/exp2f.cpp

14 ↗

(On Diff #533188)

IMO the math intrinsics should all be correctly lowered. We've traditionally expanded a lot of the f32 intrinsics in a fast math type of way. I'm fixing the ones it's reasonable to inline expand (e.g. see D153025, D153024, D153023, D153022). sqrt is also currently wrong, and I have partial patches I'm working on to fix. sqrt f32 and f64 are totally doable.

llvm.pow.f32, llvm.cos.f32 and llvm.sin.f32's correct expansions are so big I don't think it's practical to implement backend expansions for., and we should probably stop supporting the fast expansions we do for them now. The f64 versions for nearly every function (except sqrt) I think is also impractically large to handle without emitting a libcall.

[elmcdonough]

Use vendor for math functions instead of builtins.

[arsenm]

libc/src/math/gpu/frexp.cpp

19–28

See https://reviews.llvm.org/D149715, this should get a generic builtin

libc/src/math/gpu/frexpf.cpp

19

Same, https://reviews.llvm.org/D149715

Also, I think splitting the files up by FP type is pretty annoying. Many of the functions would benefit from templating

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

28

Will be correctly lowered and should use the builtin after D153025

29

Will be correctly lowered after D153024 and should use the builtin

30

In the abstract sense covered by D153024 but there's currently no defined llvm.exp10. This is asymmetrical and annoying and should be fixed, then the builtin will work

38–39

Should use the new builtin_ldexp/builtin_ldexpf

40–41

__builtin_rint

42–43

__builtin_round

45

pow/powf?

[jhuber6]

Just realized the current machinery probably doesn't handle the case where AMDGPU has a builtin but NVPTX doesn't very well right now, since we only assume one of the gpu/ or gpu/vendor entrypoints. We could probably just put the __builtin version under gpu/vendor for now and try to make it common later.

[jhuber6]

A lot of the intrinsics being implemented by @arsenm seem to be porting ocml functions to the backend. I wonder if a better approach for libc specifically would be to port those implementations to C++ so that we can share it between NVPTX and AMDGPU.

[sivachandra]

In D152603#4439280, @jhuber6 wrote:

A lot of the intrinsics being implemented by @arsenm seem to be porting ocml functions to the backend. I wonder if a better approach for libc specifically would be to port those implementations to C++ so that we can share it between NVPTX and AMDGPU.

If you can actually use the existing C++ implementations or extend/tweak the existing ones to suit GPU architecture, we would not have had to use the builtins and vendor implementations in the first place?

[elmcdonough]

Removed long double functions and replaced some vendor functions with in-progress intrinsics. pow, powf, and cosh were also added.

[arsenm]

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

26

You don't need exp_f32 or exp2_f32 anymore, the generic builtin is now correctly lowered (we're still asymmetrically missing an intrinsic for exp10). You do still need to wrap the f64 versions

35–36

No reason to wrap ldexp, just use the generic builtin

39–42

all rint and round can use the generic builtin

[jplehr]

Hi, I believe this may have broken the AMDGPU libc on GPU buildbot (that is unfortunately still in staging).
This change came in here (https://lab.llvm.org/staging/#/builders/247/builds/3734 which was broken because of another change), but the libc GPU tests are still failing (https://lab.llvm.org/staging/#/builders/247/builds/3736).

I'm probably a bit inexperienced in the whole libc infrastructure to judge this fully, but this patch would be my best guess.
Happy to help if anything needed.

[jhuber6]

Yeah this broke both buildbots, it's most likely because we're enabling some extra tests that weren't checked. Whenever a new entrypoint is added it may start building an existing test.

libc/config/gpu/entrypoints.txt

set(TARGET_LIBC_ENTRYPOINTS
     libc.src.stdio.fputs
     libc.src.stdio.stdin
     libc.src.stdio.stdout
     libc.src.stdio.stderr
 )
 
 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
     libc.src.math.roundl
 )
 
 set(TARGET_LLVMLIBC_ENTRYPOINTS
   ${TARGET_LIBC_ENTRYPOINTS}
   ${TARGET_LIBM_ENTRYPOINTS}
 )

libc/src/math/gpu/CMakeLists.txt

function(add_math_entrypoint_gpu_object name)
 
   add_entrypoint_object(
     ${name}
     ${ARGN}
   )
 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(
+  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(
   round
   SRCS
     round.cpp
   HDRS
     ../round.h
   COMPILE_OPTIONS
     -O2
 )

libc/src/math/gpu/ceil.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/ceilf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/ceill.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/copysign.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/copysignf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/copysignl.cpp

+//===-- 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

libc/src/math/gpu/fabs.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/fabsf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/fabsl.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/floor.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/floorf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/floorl.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/fma.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/fmaf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/fmax.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/fmaxf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/fmaxl.cpp

+//===-- 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

libc/src/math/gpu/fmin.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/fminf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/fminl.cpp

+//===-- 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

libc/src/math/gpu/fmod.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/fmodf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/frexp.cpp

+//===-- 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

libc/src/math/gpu/frexpf.cpp

+//===-- 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

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

 endif()
 
 # FIXME: We need a way to pass the library to only the NVTPX / AMDGPU build.
 # This shouldn't cause issues because we only link in needed symbols, but it
 # 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(
+  cos
+  SRCS
+    cos.cpp
+  HDRS
+    ../../cos.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_entrypoint_object(
+  cosf
+  SRCS
+    cosf.cpp
+  HDRS
+    ../../cosf.h
+  COMPILE_OPTIONS
+    -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(
+  exp2f
+  SRCS
+    exp2f.cpp
+  HDRS
+    ../../exp2f.h
+  COMPILE_OPTIONS
+    ${bitcode_link_flags}
+    -O2
+)
+
+add_entrypoint_object(
+  expf
+  SRCS
+    expf.cpp
+  HDRS
+    ../../expf.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(
+  ldexp
+  SRCS
+    ldexp.cpp
+  HDRS
+    ../../ldexp.h
+  COMPILE_OPTIONS
+    -O2
+)
+
+add_entrypoint_object(
+  ldexpf
+  SRCS
+    ldexpf.cpp
+  HDRS
+    ../../ldexpf.h
+  COMPILE_OPTIONS
+    -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
     sin.cpp
   HDRS
     ../../sin.h
   COMPILE_OPTIONS
     ${bitcode_link_flags}
     -O2
 )

libc/src/math/gpu/vendor/acosf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/vendor/acoshf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

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

 
 #include "declarations.h"
 #include "platform.h"
 
 #include "src/__support/macros/attributes.h"
 
 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 double cos(double x) { return __ocml_cos_f64(x); }
+LIBC_INLINE float cosf(float x) { return __ocml_cos_f32(x); }
+LIBC_INLINE float coshf(float x) { return __ocml_cosh_f32(x); }
+LIBC_INLINE float expf(float x) { return __ocml_exp_f32(x); }
+LIBC_INLINE float exp2f(float x) { return __ocml_exp2_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 double ldexp(double x, int i) { return __ocml_ldexp_f64(x, i); }
+LIBC_INLINE float ldexpf(float x, int i) { return __ocml_ldexp_f32(x, i); }
+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
 } // namespace __llvm_libc
 
 #endif // LLVM_LIBC_SRC_MATH_GPU_AMDGPU_H

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

 //===-- AMDGPU specific declarations for math support ---------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_LIBC_SRC_MATH_GPU_AMDGPU_DECLARATIONS_H
 #define LLVM_LIBC_SRC_MATH_GPU_AMDGPU_DECLARATIONS_H
 
 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_cos_f32(float);
+double __ocml_cos_f64(double);
+float __ocml_cosh_f32(float);
+float __ocml_exp_f32(float);
+float __ocml_exp2_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);
+float __ocml_ldexp_f32(float, int);
+double __ocml_ldexp_f64(double, int);
+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);
 }
 
 } // namespace __llvm_libc
 
 #endif // LLVM_LIBC_SRC_MATH_GPU_AMDGPU_DECLARATIONS_H

libc/src/math/gpu/vendor/asinf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/vendor/asinhf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/vendor/atanf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/vendor/atanhf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/vendor/cos.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(double, cos, (double x)) { return internal::cos(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H

libc/src/math/gpu/vendor/cosf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(float, cosf, (float x)) { return internal::cosf(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H

libc/src/math/gpu/vendor/coshf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/vendor/exp10f.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/vendor/exp2f.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(float, exp2f, (float x)) { return internal::exp2f(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H

libc/src/math/gpu/vendor/expf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
-}
+LLVM_LIBC_FUNCTION(float, expf, (float x)) { return internal::expf(x); }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H

libc/src/math/gpu/vendor/expm1f.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/vendor/fdim.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/vendor/fdimf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/vendor/hypot.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/vendor/hypotf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/vendor/ilogb.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/vendor/ilogbf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/vendor/ldexp.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(double, ldexp, (double x, int y)) {
+  return internal::ldexp(x, y);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H

libc/src/math/gpu/vendor/ldexpf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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"
+
+#include "common.h"
 
 namespace __llvm_libc {
 
-extern "C" {
-double __nv_sin(double);
+LLVM_LIBC_FUNCTION(float, ldexpf, (float x, int y)) {
+  return internal::ldexpf(x, y);
 }
 
 } // namespace __llvm_libc
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H

libc/src/math/gpu/vendor/llrint.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/vendor/llrintf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/vendor/llround.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

libc/src/math/gpu/vendor/llroundf.cpp

-//===-- 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.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#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

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

 //===-- NVPTX specific declarations for math support ----------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
 #define LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H
 
 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);
+double __nv_cos(double);
+float __nv_cosf(float);
+float __nv_coshf(float);
+float __nv_expf(float);
+float __nv_exp2f(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);
+double __nv_ldexp(double, int);
+float __nv_ldexpf(float, int);
+long long __nv_llrint(double);
+long long __nv_llrintf(float);
+long long __nv_llround(double);
+long long __nv_llroundf(float);
 double __nv_sin(double);
 }
 
 } // namespace __llvm_libc
 
 #endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_DECLARATIONS_H

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

 #define LLVM_LIBC_SRC_MATH_GPU_NVPTX_H
 
 #include "declarations.h"
 
 #include "src/__support/macros/attributes.h"
 
 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 double cos(double x) { return __nv_cos(x); }
+LIBC_INLINE float cosf(float x) { return __nv_cosf(x); }
+LIBC_INLINE float coshf(float x) { return __nv_coshf(x); }
+LIBC_INLINE float expf(float x) { return __nv_expf(x); }
+LIBC_INLINE float exp2f(float x) { return __nv_exp2f(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 double ldexp(double x, int i) { return __nv_ldexp(x, i); }
+LIBC_INLINE float ldexpf(float x, int i) { return __nv_ldexpf(x, i); }
+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
 } // namespace __llvm_libc
 
 #endif // LLVM_LIBC_SRC_MATH_GPU_NVPTX_H