[libc-commits] [libc] d292610 - [libc] Clean up GPU math implementations (#83133)
via libc-commits
libc-commits at lists.llvm.org
Tue Feb 27 09:39:23 PST 2024
Author: Joseph Huber
Date: 2024-02-27T11:39:13-06:00
New Revision: d29261074cd18383f4d3c84e740c66de738cb61f
URL: https://github.com/llvm/llvm-project/commit/d29261074cd18383f4d3c84e740c66de738cb61f
DIFF: https://github.com/llvm/llvm-project/commit/d29261074cd18383f4d3c84e740c66de738cb61f.diff
LOG: [libc] Clean up GPU math implementations (#83133)
Summary:
The math directory likes to do architecture specific implementations of
these math functions. For the GPU case it was complicated by the fact
that both NVPTX and AMDGPU had to go through the same code paths. Since
reworking the GPU target this is no longer the case and we can simply
use the same scheme. This patch moves all the old code into two separate
directories. This likely results in a net increase in code, but it's
easier to reason with.
Added:
libc/src/math/amdgpu/CMakeLists.txt
libc/src/math/amdgpu/acos.cpp
libc/src/math/amdgpu/acosf.cpp
libc/src/math/amdgpu/acosh.cpp
libc/src/math/amdgpu/acoshf.cpp
libc/src/math/amdgpu/asin.cpp
libc/src/math/amdgpu/asinf.cpp
libc/src/math/amdgpu/asinh.cpp
libc/src/math/amdgpu/asinhf.cpp
libc/src/math/amdgpu/atan.cpp
libc/src/math/amdgpu/atan2.cpp
libc/src/math/amdgpu/atan2f.cpp
libc/src/math/amdgpu/atanf.cpp
libc/src/math/amdgpu/atanh.cpp
libc/src/math/amdgpu/atanhf.cpp
libc/src/math/amdgpu/ceil.cpp
libc/src/math/amdgpu/ceilf.cpp
libc/src/math/amdgpu/copysign.cpp
libc/src/math/amdgpu/copysignf.cpp
libc/src/math/amdgpu/cos.cpp
libc/src/math/amdgpu/cosf.cpp
libc/src/math/amdgpu/cosh.cpp
libc/src/math/amdgpu/coshf.cpp
libc/src/math/amdgpu/declarations.h
libc/src/math/amdgpu/erf.cpp
libc/src/math/amdgpu/erff.cpp
libc/src/math/amdgpu/exp.cpp
libc/src/math/amdgpu/exp10.cpp
libc/src/math/amdgpu/exp10f.cpp
libc/src/math/amdgpu/exp2.cpp
libc/src/math/amdgpu/exp2f.cpp
libc/src/math/amdgpu/expf.cpp
libc/src/math/amdgpu/expm1.cpp
libc/src/math/amdgpu/expm1f.cpp
libc/src/math/amdgpu/fabs.cpp
libc/src/math/amdgpu/fabsf.cpp
libc/src/math/amdgpu/fdim.cpp
libc/src/math/amdgpu/fdimf.cpp
libc/src/math/amdgpu/floor.cpp
libc/src/math/amdgpu/floorf.cpp
libc/src/math/amdgpu/fma.cpp
libc/src/math/amdgpu/fmaf.cpp
libc/src/math/amdgpu/fmax.cpp
libc/src/math/amdgpu/fmaxf.cpp
libc/src/math/amdgpu/fmin.cpp
libc/src/math/amdgpu/fminf.cpp
libc/src/math/amdgpu/fmod.cpp
libc/src/math/amdgpu/fmodf.cpp
libc/src/math/amdgpu/frexp.cpp
libc/src/math/amdgpu/frexpf.cpp
libc/src/math/amdgpu/hypot.cpp
libc/src/math/amdgpu/hypotf.cpp
libc/src/math/amdgpu/ilogb.cpp
libc/src/math/amdgpu/ilogbf.cpp
libc/src/math/amdgpu/ldexp.cpp
libc/src/math/amdgpu/ldexpf.cpp
libc/src/math/amdgpu/llrint.cpp
libc/src/math/amdgpu/llrintf.cpp
libc/src/math/amdgpu/llround.cpp
libc/src/math/amdgpu/llroundf.cpp
libc/src/math/amdgpu/log.cpp
libc/src/math/amdgpu/log10.cpp
libc/src/math/amdgpu/log10f.cpp
libc/src/math/amdgpu/log1p.cpp
libc/src/math/amdgpu/log1pf.cpp
libc/src/math/amdgpu/log2.cpp
libc/src/math/amdgpu/log2f.cpp
libc/src/math/amdgpu/logb.cpp
libc/src/math/amdgpu/logbf.cpp
libc/src/math/amdgpu/logf.cpp
libc/src/math/amdgpu/lrint.cpp
libc/src/math/amdgpu/lrintf.cpp
libc/src/math/amdgpu/lround.cpp
libc/src/math/amdgpu/lroundf.cpp
libc/src/math/amdgpu/modf.cpp
libc/src/math/amdgpu/modff.cpp
libc/src/math/amdgpu/nearbyint.cpp
libc/src/math/amdgpu/nearbyintf.cpp
libc/src/math/amdgpu/nextafter.cpp
libc/src/math/amdgpu/nextafterf.cpp
libc/src/math/amdgpu/platform.h
libc/src/math/amdgpu/pow.cpp
libc/src/math/amdgpu/powf.cpp
libc/src/math/amdgpu/remainder.cpp
libc/src/math/amdgpu/remainderf.cpp
libc/src/math/amdgpu/remquo.cpp
libc/src/math/amdgpu/remquof.cpp
libc/src/math/amdgpu/rint.cpp
libc/src/math/amdgpu/rintf.cpp
libc/src/math/amdgpu/round.cpp
libc/src/math/amdgpu/roundf.cpp
libc/src/math/amdgpu/scalbn.cpp
libc/src/math/amdgpu/scalbnf.cpp
libc/src/math/amdgpu/sin.cpp
libc/src/math/amdgpu/sincos.cpp
libc/src/math/amdgpu/sincosf.cpp
libc/src/math/amdgpu/sinf.cpp
libc/src/math/amdgpu/sinh.cpp
libc/src/math/amdgpu/sinhf.cpp
libc/src/math/amdgpu/sqrt.cpp
libc/src/math/amdgpu/sqrtf.cpp
libc/src/math/amdgpu/tan.cpp
libc/src/math/amdgpu/tanf.cpp
libc/src/math/amdgpu/tanh.cpp
libc/src/math/amdgpu/tanhf.cpp
libc/src/math/amdgpu/tgamma.cpp
libc/src/math/amdgpu/tgammaf.cpp
libc/src/math/amdgpu/trunc.cpp
libc/src/math/amdgpu/truncf.cpp
libc/src/math/nvptx/CMakeLists.txt
libc/src/math/nvptx/acos.cpp
libc/src/math/nvptx/acosf.cpp
libc/src/math/nvptx/acosh.cpp
libc/src/math/nvptx/acoshf.cpp
libc/src/math/nvptx/asin.cpp
libc/src/math/nvptx/asinf.cpp
libc/src/math/nvptx/asinh.cpp
libc/src/math/nvptx/asinhf.cpp
libc/src/math/nvptx/atan.cpp
libc/src/math/nvptx/atan2.cpp
libc/src/math/nvptx/atan2f.cpp
libc/src/math/nvptx/atanf.cpp
libc/src/math/nvptx/atanh.cpp
libc/src/math/nvptx/atanhf.cpp
libc/src/math/nvptx/ceil.cpp
libc/src/math/nvptx/ceilf.cpp
libc/src/math/nvptx/copysign.cpp
libc/src/math/nvptx/copysignf.cpp
libc/src/math/nvptx/cos.cpp
libc/src/math/nvptx/cosf.cpp
libc/src/math/nvptx/cosh.cpp
libc/src/math/nvptx/coshf.cpp
libc/src/math/nvptx/declarations.h
libc/src/math/nvptx/erf.cpp
libc/src/math/nvptx/erff.cpp
libc/src/math/nvptx/exp.cpp
libc/src/math/nvptx/exp10.cpp
libc/src/math/nvptx/exp10f.cpp
libc/src/math/nvptx/exp2.cpp
libc/src/math/nvptx/exp2f.cpp
libc/src/math/nvptx/expf.cpp
libc/src/math/nvptx/expm1.cpp
libc/src/math/nvptx/expm1f.cpp
libc/src/math/nvptx/fabs.cpp
libc/src/math/nvptx/fabsf.cpp
libc/src/math/nvptx/fdim.cpp
libc/src/math/nvptx/fdimf.cpp
libc/src/math/nvptx/floor.cpp
libc/src/math/nvptx/floorf.cpp
libc/src/math/nvptx/fma.cpp
libc/src/math/nvptx/fmaf.cpp
libc/src/math/nvptx/fmax.cpp
libc/src/math/nvptx/fmaxf.cpp
libc/src/math/nvptx/fmin.cpp
libc/src/math/nvptx/fminf.cpp
libc/src/math/nvptx/fmod.cpp
libc/src/math/nvptx/fmodf.cpp
libc/src/math/nvptx/frexp.cpp
libc/src/math/nvptx/frexpf.cpp
libc/src/math/nvptx/hypot.cpp
libc/src/math/nvptx/hypotf.cpp
libc/src/math/nvptx/ilogb.cpp
libc/src/math/nvptx/ilogbf.cpp
libc/src/math/nvptx/ldexp.cpp
libc/src/math/nvptx/ldexpf.cpp
libc/src/math/nvptx/llrint.cpp
libc/src/math/nvptx/llrintf.cpp
libc/src/math/nvptx/llround.cpp
libc/src/math/nvptx/llroundf.cpp
libc/src/math/nvptx/log.cpp
libc/src/math/nvptx/log10.cpp
libc/src/math/nvptx/log10f.cpp
libc/src/math/nvptx/log1p.cpp
libc/src/math/nvptx/log1pf.cpp
libc/src/math/nvptx/log2.cpp
libc/src/math/nvptx/log2f.cpp
libc/src/math/nvptx/logb.cpp
libc/src/math/nvptx/logbf.cpp
libc/src/math/nvptx/logf.cpp
libc/src/math/nvptx/lrint.cpp
libc/src/math/nvptx/lrintf.cpp
libc/src/math/nvptx/lround.cpp
libc/src/math/nvptx/lroundf.cpp
libc/src/math/nvptx/modf.cpp
libc/src/math/nvptx/modff.cpp
libc/src/math/nvptx/nearbyint.cpp
libc/src/math/nvptx/nearbyintf.cpp
libc/src/math/nvptx/nextafter.cpp
libc/src/math/nvptx/nextafterf.cpp
libc/src/math/nvptx/nvptx.h
libc/src/math/nvptx/pow.cpp
libc/src/math/nvptx/powf.cpp
libc/src/math/nvptx/remainder.cpp
libc/src/math/nvptx/remainderf.cpp
libc/src/math/nvptx/remquo.cpp
libc/src/math/nvptx/remquof.cpp
libc/src/math/nvptx/rint.cpp
libc/src/math/nvptx/rintf.cpp
libc/src/math/nvptx/round.cpp
libc/src/math/nvptx/roundf.cpp
libc/src/math/nvptx/scalbn.cpp
libc/src/math/nvptx/scalbnf.cpp
libc/src/math/nvptx/sin.cpp
libc/src/math/nvptx/sincos.cpp
libc/src/math/nvptx/sincosf.cpp
libc/src/math/nvptx/sinf.cpp
libc/src/math/nvptx/sinh.cpp
libc/src/math/nvptx/sinhf.cpp
libc/src/math/nvptx/sqrt.cpp
libc/src/math/nvptx/sqrtf.cpp
libc/src/math/nvptx/tan.cpp
libc/src/math/nvptx/tanf.cpp
libc/src/math/nvptx/tanh.cpp
libc/src/math/nvptx/tanhf.cpp
libc/src/math/nvptx/tgamma.cpp
libc/src/math/nvptx/tgammaf.cpp
libc/src/math/nvptx/trunc.cpp
libc/src/math/nvptx/truncf.cpp
Modified:
libc/cmake/modules/LLVMLibCObjectRules.cmake
libc/src/math/CMakeLists.txt
Removed:
libc/src/math/gpu/CMakeLists.txt
libc/src/math/gpu/ceil.cpp
libc/src/math/gpu/ceilf.cpp
libc/src/math/gpu/copysign.cpp
libc/src/math/gpu/copysignf.cpp
libc/src/math/gpu/fabs.cpp
libc/src/math/gpu/fabsf.cpp
libc/src/math/gpu/floor.cpp
libc/src/math/gpu/floorf.cpp
libc/src/math/gpu/fma.cpp
libc/src/math/gpu/fmaf.cpp
libc/src/math/gpu/fmax.cpp
libc/src/math/gpu/fmaxf.cpp
libc/src/math/gpu/fmin.cpp
libc/src/math/gpu/fminf.cpp
libc/src/math/gpu/fmod.cpp
libc/src/math/gpu/fmodf.cpp
libc/src/math/gpu/llround.cpp
libc/src/math/gpu/llroundf.cpp
libc/src/math/gpu/lround.cpp
libc/src/math/gpu/lroundf.cpp
libc/src/math/gpu/modf.cpp
libc/src/math/gpu/modff.cpp
libc/src/math/gpu/nearbyint.cpp
libc/src/math/gpu/nearbyintf.cpp
libc/src/math/gpu/remainder.cpp
libc/src/math/gpu/remainderf.cpp
libc/src/math/gpu/rint.cpp
libc/src/math/gpu/rintf.cpp
libc/src/math/gpu/round.cpp
libc/src/math/gpu/roundf.cpp
libc/src/math/gpu/sinh.cpp
libc/src/math/gpu/sqrt.cpp
libc/src/math/gpu/sqrtf.cpp
libc/src/math/gpu/tan.cpp
libc/src/math/gpu/tanh.cpp
libc/src/math/gpu/trunc.cpp
libc/src/math/gpu/truncf.cpp
libc/src/math/gpu/vendor/CMakeLists.txt
libc/src/math/gpu/vendor/acos.cpp
libc/src/math/gpu/vendor/acosf.cpp
libc/src/math/gpu/vendor/acosh.cpp
libc/src/math/gpu/vendor/acoshf.cpp
libc/src/math/gpu/vendor/amdgpu/amdgpu.h
libc/src/math/gpu/vendor/amdgpu/declarations.h
libc/src/math/gpu/vendor/amdgpu/platform.h
libc/src/math/gpu/vendor/asin.cpp
libc/src/math/gpu/vendor/asinf.cpp
libc/src/math/gpu/vendor/asinh.cpp
libc/src/math/gpu/vendor/asinhf.cpp
libc/src/math/gpu/vendor/atan.cpp
libc/src/math/gpu/vendor/atan2.cpp
libc/src/math/gpu/vendor/atan2f.cpp
libc/src/math/gpu/vendor/atanf.cpp
libc/src/math/gpu/vendor/atanh.cpp
libc/src/math/gpu/vendor/atanhf.cpp
libc/src/math/gpu/vendor/common.h
libc/src/math/gpu/vendor/cos.cpp
libc/src/math/gpu/vendor/cosf.cpp
libc/src/math/gpu/vendor/cosh.cpp
libc/src/math/gpu/vendor/coshf.cpp
libc/src/math/gpu/vendor/erf.cpp
libc/src/math/gpu/vendor/erff.cpp
libc/src/math/gpu/vendor/exp.cpp
libc/src/math/gpu/vendor/exp10.cpp
libc/src/math/gpu/vendor/exp10f.cpp
libc/src/math/gpu/vendor/exp2.cpp
libc/src/math/gpu/vendor/exp2f.cpp
libc/src/math/gpu/vendor/expf.cpp
libc/src/math/gpu/vendor/expm1.cpp
libc/src/math/gpu/vendor/expm1f.cpp
libc/src/math/gpu/vendor/fdim.cpp
libc/src/math/gpu/vendor/fdimf.cpp
libc/src/math/gpu/vendor/frexp.cpp
libc/src/math/gpu/vendor/frexpf.cpp
libc/src/math/gpu/vendor/hypot.cpp
libc/src/math/gpu/vendor/hypotf.cpp
libc/src/math/gpu/vendor/ilogb.cpp
libc/src/math/gpu/vendor/ilogbf.cpp
libc/src/math/gpu/vendor/ldexp.cpp
libc/src/math/gpu/vendor/ldexpf.cpp
libc/src/math/gpu/vendor/llrint.cpp
libc/src/math/gpu/vendor/llrintf.cpp
libc/src/math/gpu/vendor/log.cpp
libc/src/math/gpu/vendor/log10.cpp
libc/src/math/gpu/vendor/log10f.cpp
libc/src/math/gpu/vendor/log1p.cpp
libc/src/math/gpu/vendor/log1pf.cpp
libc/src/math/gpu/vendor/log2.cpp
libc/src/math/gpu/vendor/log2f.cpp
libc/src/math/gpu/vendor/logb.cpp
libc/src/math/gpu/vendor/logbf.cpp
libc/src/math/gpu/vendor/logf.cpp
libc/src/math/gpu/vendor/lrint.cpp
libc/src/math/gpu/vendor/lrintf.cpp
libc/src/math/gpu/vendor/nextafter.cpp
libc/src/math/gpu/vendor/nextafterf.cpp
libc/src/math/gpu/vendor/nvptx/declarations.h
libc/src/math/gpu/vendor/nvptx/nvptx.h
libc/src/math/gpu/vendor/pow.cpp
libc/src/math/gpu/vendor/powf.cpp
libc/src/math/gpu/vendor/remquo.cpp
libc/src/math/gpu/vendor/remquof.cpp
libc/src/math/gpu/vendor/scalbn.cpp
libc/src/math/gpu/vendor/scalbnf.cpp
libc/src/math/gpu/vendor/sin.cpp
libc/src/math/gpu/vendor/sincos.cpp
libc/src/math/gpu/vendor/sincosf.cpp
libc/src/math/gpu/vendor/sinf.cpp
libc/src/math/gpu/vendor/sinh.cpp
libc/src/math/gpu/vendor/sinhf.cpp
libc/src/math/gpu/vendor/tan.cpp
libc/src/math/gpu/vendor/tanf.cpp
libc/src/math/gpu/vendor/tanh.cpp
libc/src/math/gpu/vendor/tanhf.cpp
libc/src/math/gpu/vendor/tgamma.cpp
libc/src/math/gpu/vendor/tgammaf.cpp
################################################################################
diff --git a/libc/cmake/modules/LLVMLibCObjectRules.cmake b/libc/cmake/modules/LLVMLibCObjectRules.cmake
index 8a84c82206ba6c..0649e9f7a76709 100644
--- a/libc/cmake/modules/LLVMLibCObjectRules.cmake
+++ b/libc/cmake/modules/LLVMLibCObjectRules.cmake
@@ -307,7 +307,7 @@ function(create_entrypoint_object fq_target_name)
${fq_target_name}
PROPERTIES
ENTRYPOINT_NAME ${ADD_ENTRYPOINT_OBJ_NAME}
- TARGET_TYPE ${ENTRYPOINT_OBJ_TARGET_TYPE}
+ TARGET_TYPE ${entrypoint_target_type}
OBJECT_FILE "$<TARGET_OBJECTS:${fq_target_name}>"
CXX_STANDARD ${ADD_ENTRYPOINT_OBJ_CXX_STANDARD}
DEPS "${fq_deps_list}"
diff --git a/libc/src/math/CMakeLists.txt b/libc/src/math/CMakeLists.txt
index efa14c42e389b3..882befd9f7e7ff 100644
--- a/libc/src/math/CMakeLists.txt
+++ b/libc/src/math/CMakeLists.txt
@@ -1,9 +1,6 @@
add_subdirectory(generic)
if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/${LIBC_TARGET_ARCHITECTURE})
add_subdirectory(${LIBC_TARGET_ARCHITECTURE})
-elseif(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/${LIBC_TARGET_OS})
- # TODO: We should split this into 'nvptx' and 'amdgpu' for the GPU build.
- add_subdirectory(${LIBC_TARGET_OS})
endif()
function(add_math_entrypoint_object name)
@@ -11,7 +8,6 @@ function(add_math_entrypoint_object name)
# that first and return early if we are able to add an alias target for the
# machine specific implementation.
get_fq_target_name("${LIBC_TARGET_ARCHITECTURE}.${name}" fq_machine_specific_target_name)
- get_fq_target_name("${LIBC_TARGET_OS}.${name}" fq_os_specific_target_name)
if(TARGET ${fq_machine_specific_target_name})
add_entrypoint_object(
${name}
@@ -20,28 +16,6 @@ function(add_math_entrypoint_object name)
.${LIBC_TARGET_ARCHITECTURE}.${name}
)
return()
- elseif(TARGET ${fq_os_specific_target_name})
- add_entrypoint_object(
- ${name}
- ALIAS
- DEPENDS
- .${LIBC_TARGET_OS}.${name}
- )
- return()
- endif()
-
- # The GPU optionally depends on vendor libraries. If we emitted one of these
- # entrypoints it means the user requested it and we should use it instead.
- get_fq_target_name("${LIBC_TARGET_OS}.vendor.${name}" fq_vendor_specific_target_name)
- if(TARGET ${fq_vendor_specific_target_name})
- add_entrypoint_object(
- ${name}
- ALIAS
- DEPENDS
- .${LIBC_TARGET_OS}.vendor.${name}
- VENDOR
- )
- return()
endif()
get_fq_target_name("generic.${name}" fq_generic_target_name)
diff --git a/libc/src/math/gpu/vendor/CMakeLists.txt b/libc/src/math/amdgpu/CMakeLists.txt
similarity index 59%
rename from libc/src/math/gpu/vendor/CMakeLists.txt
rename to libc/src/math/amdgpu/CMakeLists.txt
index 36087ade63bfcd..cb77341aa50522 100644
--- a/libc/src/math/gpu/vendor/CMakeLists.txt
+++ b/libc/src/math/amdgpu/CMakeLists.txt
@@ -1,39 +1,360 @@
+# Math functions not yet available in the libc project, or those not yet tuned
+# for GPU workloads are provided as wrappers over vendor libraries. If we find
+# them ahead of time we will import them statically. Otherwise, we will keep
+# them as external references and expect them to be resolved by the user when
+# they compile. In the future,we will use implementations from the 'libc'
+# project and not provide these wrappers.
find_package(AMDDeviceLibs QUIET HINTS ${CMAKE_INSTALL_PREFIX} PATHS /opt/rocm)
if(AMDDeviceLibs_FOUND)
message(STATUS "Found the ROCm device library. Implementations falling back "
"to the vendor libraries will be resolved statically.")
get_target_property(ocml_path ocml IMPORTED_LOCATION)
- list(APPEND bitcode_link_flags
- "SHELL:-Xclang -mlink-builtin-bitcode -Xclang ${ocml_path}")
+ set(bitcode_link_flags
+ "SHELL:-Xclang -mlink-builtin-bitcode -Xclang ${ocml_path}")
else()
message(STATUS "Could not find the ROCm device library. Unimplemented "
"functions will be an external reference to the vendor libraries.")
endif()
-if(CUDAToolkit_FOUND)
- set(libdevice_path ${CUDAToolkit_BIN_DIR}/../nvvm/libdevice/libdevice.10.bc)
- if (EXISTS ${libdevice_path})
- message(STATUS "Found the CUDA device library. Implementations falling back "
- "to the vendor libraries will be resolved statically.")
- list(APPEND bitcode_link_flags
- "SHELL:-Xclang -mlink-builtin-bitcode -Xclang ${libdevice_path}")
- endif()
-else()
- message(STATUS "Could not find the CUDA device library. Unimplemented "
- "functions will be an external reference to the vendor libraries.")
-endif()
+add_entrypoint_object(
+ ceil
+ SRCS
+ ceil.cpp
+ HDRS
+ ../ceil.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ ceilf
+ SRCS
+ ceilf.cpp
+ HDRS
+ ../ceilf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ copysign
+ SRCS
+ copysign.cpp
+ HDRS
+ ../copysign.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ copysignf
+ SRCS
+ copysignf.cpp
+ HDRS
+ ../copysignf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fabs
+ SRCS
+ fabs.cpp
+ HDRS
+ ../fabs.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fabsf
+ SRCS
+ fabsf.cpp
+ HDRS
+ ../fabsf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ floor
+ SRCS
+ floor.cpp
+ HDRS
+ ../floor.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ floorf
+ SRCS
+ floorf.cpp
+ HDRS
+ ../floorf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fma
+ SRCS
+ fma.cpp
+ HDRS
+ ../fma.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fmaf
+ SRCS
+ fmaf.cpp
+ HDRS
+ ../fmaf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fmax
+ SRCS
+ fmax.cpp
+ HDRS
+ ../fmax.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fmaxf
+ SRCS
+ fmaxf.cpp
+ HDRS
+ ../fmaxf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fmin
+ SRCS
+ fmin.cpp
+ HDRS
+ ../fmin.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fminf
+ SRCS
+ fminf.cpp
+ HDRS
+ ../fminf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fmod
+ SRCS
+ fmod.cpp
+ HDRS
+ ../fmod.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fmodf
+ SRCS
+ fmodf.cpp
+ HDRS
+ ../fmodf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ lround
+ SRCS
+ lround.cpp
+ HDRS
+ ../lround.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ lroundf
+ SRCS
+ lroundf.cpp
+ HDRS
+ ../lroundf.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(
+ modf
+ SRCS
+ modf.cpp
+ HDRS
+ ../modf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ modff
+ SRCS
+ modff.cpp
+ HDRS
+ ../modff.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ nearbyint
+ SRCS
+ nearbyint.cpp
+ HDRS
+ ../nearbyint.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ nearbyintf
+ SRCS
+ nearbyintf.cpp
+ HDRS
+ ../nearbyintf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ remainder
+ SRCS
+ remainder.cpp
+ HDRS
+ ../remainder.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ remainderf
+ SRCS
+ remainderf.cpp
+ HDRS
+ ../remainderf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ rint
+ SRCS
+ rint.cpp
+ HDRS
+ ../rint.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ rintf
+ SRCS
+ rintf.cpp
+ HDRS
+ ../rintf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ round
+ SRCS
+ round.cpp
+ HDRS
+ ../round.h
+ COMPILE_OPTIONS
+ -O2
+)
-# 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(
+ sqrt
+ SRCS
+ sqrt.cpp
+ HDRS
+ ../sqrt.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ sqrtf
+ SRCS
+ sqrtf.cpp
+ HDRS
+ ../sqrtf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ trunc
+ SRCS
+ trunc.cpp
+ HDRS
+ ../trunc.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ truncf
+ SRCS
+ truncf.cpp
+ HDRS
+ ../truncf.h
+ COMPILE_OPTIONS
+ -O2
+)
+# The following functions currently are not implemented natively and borrow from
+# existing implementations. This will be removed in the future.
add_entrypoint_object(
acos
SRCS
acos.cpp
HDRS
- ../../acos.h
+ ../acos.h
+ VENDOR
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
@@ -44,10 +365,11 @@ add_entrypoint_object(
SRCS
acosf.cpp
HDRS
- ../../acosf.h
+ ../acosf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -55,10 +377,11 @@ add_entrypoint_object(
SRCS
acosh.cpp
HDRS
- ../../acosh.h
+ ../acosh.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -66,10 +389,11 @@ add_entrypoint_object(
SRCS
acoshf.cpp
HDRS
- ../../acoshf.h
+ ../acoshf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -77,10 +401,11 @@ add_entrypoint_object(
SRCS
asin.cpp
HDRS
- ../../asin.h
+ ../asin.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -88,10 +413,11 @@ add_entrypoint_object(
SRCS
asinf.cpp
HDRS
- ../../asinf.h
+ ../asinf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -99,10 +425,11 @@ add_entrypoint_object(
SRCS
asinh.cpp
HDRS
- ../../asinh.h
+ ../asinh.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -110,10 +437,11 @@ add_entrypoint_object(
SRCS
atan.cpp
HDRS
- ../../atan.h
+ ../atan.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -121,10 +449,11 @@ add_entrypoint_object(
SRCS
atanf.cpp
HDRS
- ../../atanf.h
+ ../atanf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -132,10 +461,11 @@ add_entrypoint_object(
SRCS
atan2.cpp
HDRS
- ../../atan2.h
+ ../atan2.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -143,10 +473,11 @@ add_entrypoint_object(
SRCS
atan2f.cpp
HDRS
- ../../atan2f.h
+ ../atan2f.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -154,10 +485,11 @@ add_entrypoint_object(
SRCS
atanh.cpp
HDRS
- ../../atanh.h
+ ../atanh.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -165,10 +497,11 @@ add_entrypoint_object(
SRCS
atanhf.cpp
HDRS
- ../../atanhf.h
+ ../atanhf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -176,10 +509,11 @@ add_entrypoint_object(
SRCS
cos.cpp
HDRS
- ../../cos.h
+ ../cos.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -187,10 +521,11 @@ add_entrypoint_object(
SRCS
cosf.cpp
HDRS
- ../../cosf.h
+ ../cosf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -198,10 +533,11 @@ add_entrypoint_object(
SRCS
cosh.cpp
HDRS
- ../../cosh.h
+ ../cosh.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -209,10 +545,11 @@ add_entrypoint_object(
SRCS
coshf.cpp
HDRS
- ../../coshf.h
+ ../coshf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -220,10 +557,11 @@ add_entrypoint_object(
SRCS
erf.cpp
HDRS
- ../../erf.h
+ ../erf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -231,10 +569,11 @@ add_entrypoint_object(
SRCS
erff.cpp
HDRS
- ../../erff.h
+ ../erff.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -242,10 +581,11 @@ add_entrypoint_object(
SRCS
exp.cpp
HDRS
- ../../exp.h
+ ../exp.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -253,10 +593,11 @@ add_entrypoint_object(
SRCS
exp10.cpp
HDRS
- ../../exp10.h
+ ../exp10.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -264,10 +605,11 @@ add_entrypoint_object(
SRCS
exp10f.cpp
HDRS
- ../../exp10f.h
+ ../exp10f.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -275,10 +617,11 @@ add_entrypoint_object(
SRCS
exp2.cpp
HDRS
- ../../exp2.h
+ ../exp2.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -286,10 +629,11 @@ add_entrypoint_object(
SRCS
exp2f.cpp
HDRS
- ../../exp2f.h
+ ../exp2f.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -297,10 +641,11 @@ add_entrypoint_object(
SRCS
expf.cpp
HDRS
- ../../expf.h
+ ../expf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -308,10 +653,11 @@ add_entrypoint_object(
SRCS
expm1.cpp
HDRS
- ../../expm1.h
+ ../expm1.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -319,10 +665,11 @@ add_entrypoint_object(
SRCS
expm1f.cpp
HDRS
- ../../expm1f.h
+ ../expm1f.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -330,10 +677,11 @@ add_entrypoint_object(
SRCS
fdim.cpp
HDRS
- ../../fdim.h
+ ../fdim.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -341,10 +689,11 @@ add_entrypoint_object(
SRCS
fdimf.cpp
HDRS
- ../../fdimf.h
+ ../fdimf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -352,10 +701,11 @@ add_entrypoint_object(
SRCS
hypot.cpp
HDRS
- ../../hypot.h
+ ../hypot.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -363,10 +713,11 @@ add_entrypoint_object(
SRCS
hypotf.cpp
HDRS
- ../../hypotf.h
+ ../hypotf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -374,10 +725,11 @@ add_entrypoint_object(
SRCS
ilogb.cpp
HDRS
- ../../ilogb.h
+ ../ilogb.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -385,10 +737,11 @@ add_entrypoint_object(
SRCS
ilogbf.cpp
HDRS
- ../../ilogbf.h
+ ../ilogbf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -396,10 +749,11 @@ add_entrypoint_object(
SRCS
log10.cpp
HDRS
- ../../log10.h
+ ../log10.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -407,10 +761,11 @@ add_entrypoint_object(
SRCS
log10f.cpp
HDRS
- ../../log10f.h
+ ../log10f.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -418,10 +773,11 @@ add_entrypoint_object(
SRCS
log2.cpp
HDRS
- ../../log2.h
+ ../log2.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -429,10 +785,11 @@ add_entrypoint_object(
SRCS
log2f.cpp
HDRS
- ../../log2f.h
+ ../log2f.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -440,10 +797,11 @@ add_entrypoint_object(
SRCS
log.cpp
HDRS
- ../../log.h
+ ../log.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -451,10 +809,11 @@ add_entrypoint_object(
SRCS
logf.cpp
HDRS
- ../../logf.h
+ ../logf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -462,10 +821,11 @@ add_entrypoint_object(
SRCS
lrint.cpp
HDRS
- ../../lrint.h
+ ../lrint.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -473,10 +833,11 @@ add_entrypoint_object(
SRCS
lrintf.cpp
HDRS
- ../../lrintf.h
+ ../lrintf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -484,10 +845,11 @@ add_entrypoint_object(
SRCS
ldexp.cpp
HDRS
- ../../ldexp.h
+ ../ldexp.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -495,10 +857,11 @@ add_entrypoint_object(
SRCS
ldexpf.cpp
HDRS
- ../../ldexpf.h
+ ../ldexpf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -506,10 +869,11 @@ add_entrypoint_object(
SRCS
log1p.cpp
HDRS
- ../../log1p.h
+ ../log1p.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -517,10 +881,11 @@ add_entrypoint_object(
SRCS
log1pf.cpp
HDRS
- ../../log1pf.h
+ ../log1pf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -528,10 +893,11 @@ add_entrypoint_object(
SRCS
llrint.cpp
HDRS
- ../../llrint.h
+ ../llrint.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -539,10 +905,11 @@ add_entrypoint_object(
SRCS
llrintf.cpp
HDRS
- ../../llrintf.h
+ ../llrintf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -550,10 +917,11 @@ add_entrypoint_object(
SRCS
remquo.cpp
HDRS
- ../../remquo.h
+ ../remquo.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -561,10 +929,11 @@ add_entrypoint_object(
SRCS
remquof.cpp
HDRS
- ../../remquof.h
+ ../remquof.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -572,10 +941,11 @@ add_entrypoint_object(
SRCS
scalbn.cpp
HDRS
- ../../scalbn.h
+ ../scalbn.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -583,10 +953,11 @@ add_entrypoint_object(
SRCS
scalbnf.cpp
HDRS
- ../../scalbnf.h
+ ../scalbnf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
@@ -595,10 +966,11 @@ add_entrypoint_object(
SRCS
nextafter.cpp
HDRS
- ../../nextafter.h
+ ../nextafter.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -606,10 +978,11 @@ add_entrypoint_object(
SRCS
nextafterf.cpp
HDRS
- ../../nextafterf.h
+ ../nextafterf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -617,10 +990,11 @@ add_entrypoint_object(
SRCS
pow.cpp
HDRS
- ../../pow.h
+ ../pow.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -628,10 +1002,11 @@ add_entrypoint_object(
SRCS
powf.cpp
HDRS
- ../../powf.h
+ ../powf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -639,10 +1014,11 @@ add_entrypoint_object(
SRCS
sin.cpp
HDRS
- ../../sin.h
+ ../sin.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -650,10 +1026,11 @@ add_entrypoint_object(
SRCS
sinf.cpp
HDRS
- ../../sinf.h
+ ../sinf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -661,10 +1038,11 @@ add_entrypoint_object(
SRCS
sincos.cpp
HDRS
- ../../sincos.h
+ ../sincos.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -672,10 +1050,11 @@ add_entrypoint_object(
SRCS
sincosf.cpp
HDRS
- ../../sincosf.h
+ ../sincosf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -683,10 +1062,11 @@ add_entrypoint_object(
SRCS
sinh.cpp
HDRS
- ../../sinh.h
+ ../sinh.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -694,10 +1074,11 @@ add_entrypoint_object(
SRCS
sinhf.cpp
HDRS
- ../../sinhf.h
+ ../sinhf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -705,10 +1086,11 @@ add_entrypoint_object(
SRCS
tan.cpp
HDRS
- ../../tan.h
+ ../tan.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -716,10 +1098,11 @@ add_entrypoint_object(
SRCS
tanf.cpp
HDRS
- ../../tanf.h
+ ../tanf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -727,10 +1110,11 @@ add_entrypoint_object(
SRCS
tanh.cpp
HDRS
- ../../tanh.h
+ ../tanh.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -738,10 +1122,11 @@ add_entrypoint_object(
SRCS
tanhf.cpp
HDRS
- ../../tanhf.h
+ ../tanhf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -749,10 +1134,11 @@ add_entrypoint_object(
SRCS
tgamma.cpp
HDRS
- ../../tgamma.h
+ ../tgamma.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -760,10 +1146,11 @@ add_entrypoint_object(
SRCS
tgammaf.cpp
HDRS
- ../../tgammaf.h
+ ../tgammaf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -771,10 +1158,11 @@ add_entrypoint_object(
SRCS
frexp.cpp
HDRS
- ../../frexp.h
+ ../frexp.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
add_entrypoint_object(
@@ -782,8 +1170,9 @@ add_entrypoint_object(
SRCS
frexpf.cpp
HDRS
- ../../frexpf.h
+ ../frexpf.h
COMPILE_OPTIONS
${bitcode_link_flags}
-O2
+ VENDOR
)
diff --git a/libc/src/math/amdgpu/acos.cpp b/libc/src/math/amdgpu/acos.cpp
new file mode 100644
index 00000000000000..b1e30fef82ded9
--- /dev/null
+++ b/libc/src/math/amdgpu/acos.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU acos function ---------------------------===//
+//
+// 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/acos.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, acos, (double x)) { return __ocml_acos_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/acosf.cpp b/libc/src/math/amdgpu/acosf.cpp
new file mode 100644
index 00000000000000..4c2dd4bcf4353e
--- /dev/null
+++ b/libc/src/math/amdgpu/acosf.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/acosf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, acosf, (float x)) { return __ocml_acos_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/acosh.cpp b/libc/src/math/amdgpu/acosh.cpp
new file mode 100644
index 00000000000000..dcdeeab29454e8
--- /dev/null
+++ b/libc/src/math/amdgpu/acosh.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU acosh function --------------------------===//
+//
+// 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/acosh.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, acosh, (double x)) { return __ocml_acosh_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/acoshf.cpp b/libc/src/math/amdgpu/acoshf.cpp
new file mode 100644
index 00000000000000..52baa2eaecc799
--- /dev/null
+++ b/libc/src/math/amdgpu/acoshf.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/acoshf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, acoshf, (float x)) { return __ocml_acosh_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/asin.cpp b/libc/src/math/amdgpu/asin.cpp
new file mode 100644
index 00000000000000..835c317112e2e2
--- /dev/null
+++ b/libc/src/math/amdgpu/asin.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU asin function ---------------------------===//
+//
+// 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/asin.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, asin, (double x)) { return __ocml_asin_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/asinf.cpp b/libc/src/math/amdgpu/asinf.cpp
new file mode 100644
index 00000000000000..72c45d5cf1723d
--- /dev/null
+++ b/libc/src/math/amdgpu/asinf.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/asinf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, asinf, (float x)) { return __ocml_asin_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/asinh.cpp b/libc/src/math/amdgpu/asinh.cpp
new file mode 100644
index 00000000000000..7a9f7ea4e98871
--- /dev/null
+++ b/libc/src/math/amdgpu/asinh.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU asinh function --------------------------===//
+//
+// 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/asinh.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, asinh, (double x)) { return __ocml_asinh_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/asinhf.cpp b/libc/src/math/amdgpu/asinhf.cpp
new file mode 100644
index 00000000000000..28d6bde5c91841
--- /dev/null
+++ b/libc/src/math/amdgpu/asinhf.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/asinhf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, asinhf, (float x)) { return __ocml_asinh_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/atan.cpp b/libc/src/math/amdgpu/atan.cpp
new file mode 100644
index 00000000000000..a1fa38ba451dee
--- /dev/null
+++ b/libc/src/math/amdgpu/atan.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU atan function ---------------------------===//
+//
+// 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/atan.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, atan, (double x)) { return __ocml_atan_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/atan2.cpp b/libc/src/math/amdgpu/atan2.cpp
new file mode 100644
index 00000000000000..9cfdba75eb8d61
--- /dev/null
+++ b/libc/src/math/amdgpu/atan2.cpp
@@ -0,0 +1,20 @@
+//===-- Implementation of the GPU atan2 function --------------------------===//
+//
+// 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/atan2.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, atan2, (double x, double y)) {
+ return __ocml_atan2_f64(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/atan2f.cpp b/libc/src/math/amdgpu/atan2f.cpp
new file mode 100644
index 00000000000000..ef56293b4caf53
--- /dev/null
+++ b/libc/src/math/amdgpu/atan2f.cpp
@@ -0,0 +1,20 @@
+//===-- Implementation of the GPU atan2f function -------------------------===//
+//
+// 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/atan2f.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, atan2f, (float x, float y)) {
+ return __ocml_atan2_f32(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/atanf.cpp b/libc/src/math/amdgpu/atanf.cpp
new file mode 100644
index 00000000000000..bbcceca3ed092d
--- /dev/null
+++ b/libc/src/math/amdgpu/atanf.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/atanf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, atanf, (float x)) { return __ocml_atan_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/atanh.cpp b/libc/src/math/amdgpu/atanh.cpp
new file mode 100644
index 00000000000000..ec462586450bf6
--- /dev/null
+++ b/libc/src/math/amdgpu/atanh.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU atanh function --------------------------===//
+//
+// 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/atanh.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, atanh, (double x)) { return __ocml_atanh_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/atanhf.cpp b/libc/src/math/amdgpu/atanhf.cpp
new file mode 100644
index 00000000000000..227269369ab6bd
--- /dev/null
+++ b/libc/src/math/amdgpu/atanhf.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/atanhf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, atanhf, (float x)) { return __ocml_atanh_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/ceil.cpp b/libc/src/math/amdgpu/ceil.cpp
similarity index 100%
rename from libc/src/math/gpu/ceil.cpp
rename to libc/src/math/amdgpu/ceil.cpp
diff --git a/libc/src/math/gpu/ceilf.cpp b/libc/src/math/amdgpu/ceilf.cpp
similarity index 100%
rename from libc/src/math/gpu/ceilf.cpp
rename to libc/src/math/amdgpu/ceilf.cpp
diff --git a/libc/src/math/gpu/copysign.cpp b/libc/src/math/amdgpu/copysign.cpp
similarity index 100%
rename from libc/src/math/gpu/copysign.cpp
rename to libc/src/math/amdgpu/copysign.cpp
diff --git a/libc/src/math/gpu/copysignf.cpp b/libc/src/math/amdgpu/copysignf.cpp
similarity index 100%
rename from libc/src/math/gpu/copysignf.cpp
rename to libc/src/math/amdgpu/copysignf.cpp
diff --git a/libc/src/math/amdgpu/cos.cpp b/libc/src/math/amdgpu/cos.cpp
new file mode 100644
index 00000000000000..68239d93378014
--- /dev/null
+++ b/libc/src/math/amdgpu/cos.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/cos.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, cos, (double x)) { return __ocml_cos_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/cosf.cpp b/libc/src/math/amdgpu/cosf.cpp
new file mode 100644
index 00000000000000..a60e9ea28907c1
--- /dev/null
+++ b/libc/src/math/amdgpu/cosf.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/cosf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, cosf, (float x)) { return __ocml_cos_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/cosh.cpp b/libc/src/math/amdgpu/cosh.cpp
new file mode 100644
index 00000000000000..b71df0c0170c3a
--- /dev/null
+++ b/libc/src/math/amdgpu/cosh.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the cosh 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/cosh.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, cosh, (double x)) { return __ocml_cosh_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/coshf.cpp b/libc/src/math/amdgpu/coshf.cpp
new file mode 100644
index 00000000000000..699fb0478aee85
--- /dev/null
+++ b/libc/src/math/amdgpu/coshf.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/coshf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, coshf, (float x)) { return __ocml_cosh_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/amdgpu/declarations.h b/libc/src/math/amdgpu/declarations.h
similarity index 99%
rename from libc/src/math/gpu/vendor/amdgpu/declarations.h
rename to libc/src/math/amdgpu/declarations.h
index 7a01cbc6ab19d0..780d5f0a114032 100644
--- a/libc/src/math/gpu/vendor/amdgpu/declarations.h
+++ b/libc/src/math/amdgpu/declarations.h
@@ -9,6 +9,8 @@
#ifndef LLVM_LIBC_SRC_MATH_GPU_AMDGPU_DECLARATIONS_H
#define LLVM_LIBC_SRC_MATH_GPU_AMDGPU_DECLARATIONS_H
+#include "platform.h"
+
#include "src/__support/GPU/utils.h"
namespace LIBC_NAMESPACE {
diff --git a/libc/src/math/amdgpu/erf.cpp b/libc/src/math/amdgpu/erf.cpp
new file mode 100644
index 00000000000000..7a464550c7e061
--- /dev/null
+++ b/libc/src/math/amdgpu/erf.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU erf function ----------------------------===//
+//
+// 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/erf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, erf, (double x)) { return __ocml_erf_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/erff.cpp b/libc/src/math/amdgpu/erff.cpp
new file mode 100644
index 00000000000000..1f77d08585a3d4
--- /dev/null
+++ b/libc/src/math/amdgpu/erff.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU erff function ---------------------------===//
+//
+// 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/erff.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, erff, (float x)) { return __ocml_erf_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/exp.cpp b/libc/src/math/amdgpu/exp.cpp
similarity index 84%
rename from libc/src/math/gpu/vendor/exp.cpp
rename to libc/src/math/amdgpu/exp.cpp
index ee5a22019f6a54..8590ac75901930 100644
--- a/libc/src/math/gpu/vendor/exp.cpp
+++ b/libc/src/math/amdgpu/exp.cpp
@@ -9,10 +9,10 @@
#include "src/math/exp.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, exp, (double x)) { return internal::exp(x); }
+LLVM_LIBC_FUNCTION(double, exp, (double x)) { return __builtin_exp(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/exp10.cpp b/libc/src/math/amdgpu/exp10.cpp
new file mode 100644
index 00000000000000..17d8f3350ac24b
--- /dev/null
+++ b/libc/src/math/amdgpu/exp10.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU exp10 function --------------------------===//
+//
+// 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/exp10.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, exp10, (double x)) { return __ocml_exp10_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/exp10f.cpp b/libc/src/math/amdgpu/exp10f.cpp
new file mode 100644
index 00000000000000..ddab555a8fbd30
--- /dev/null
+++ b/libc/src/math/amdgpu/exp10f.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/exp10f.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, exp10f, (float x)) { return __ocml_exp10_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/exp2.cpp b/libc/src/math/amdgpu/exp2.cpp
new file mode 100644
index 00000000000000..dfbb1f80d12949
--- /dev/null
+++ b/libc/src/math/amdgpu/exp2.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU exp2 function ---------------------------===//
+//
+// 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/exp2.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, exp2, (double x)) { return __ocml_exp2_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/exp2f.cpp b/libc/src/math/amdgpu/exp2f.cpp
new file mode 100644
index 00000000000000..016dfe3a622237
--- /dev/null
+++ b/libc/src/math/amdgpu/exp2f.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/exp2f.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, exp2f, (float x)) { return __ocml_exp2_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/expf.cpp b/libc/src/math/amdgpu/expf.cpp
similarity index 84%
rename from libc/src/math/gpu/vendor/expf.cpp
rename to libc/src/math/amdgpu/expf.cpp
index 89c194e4bc291e..d682f6293a6cb7 100644
--- a/libc/src/math/gpu/vendor/expf.cpp
+++ b/libc/src/math/amdgpu/expf.cpp
@@ -9,10 +9,10 @@
#include "src/math/expf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, expf, (float x)) { return internal::expf(x); }
+LLVM_LIBC_FUNCTION(float, expf, (float x)) { return __builtin_expf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/expm1.cpp b/libc/src/math/amdgpu/expm1.cpp
new file mode 100644
index 00000000000000..d2ac28ae6a3e29
--- /dev/null
+++ b/libc/src/math/amdgpu/expm1.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU expm1 function --------------------------===//
+//
+// 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/expm1.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, expm1, (double x)) { return __ocml_expm1_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/expm1f.cpp b/libc/src/math/amdgpu/expm1f.cpp
new file mode 100644
index 00000000000000..0ffe1a362af12a
--- /dev/null
+++ b/libc/src/math/amdgpu/expm1f.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/expm1f.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, expm1f, (float x)) { return __ocml_expm1_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/fabs.cpp b/libc/src/math/amdgpu/fabs.cpp
similarity index 100%
rename from libc/src/math/gpu/fabs.cpp
rename to libc/src/math/amdgpu/fabs.cpp
diff --git a/libc/src/math/gpu/fabsf.cpp b/libc/src/math/amdgpu/fabsf.cpp
similarity index 100%
rename from libc/src/math/gpu/fabsf.cpp
rename to libc/src/math/amdgpu/fabsf.cpp
diff --git a/libc/src/math/amdgpu/fdim.cpp b/libc/src/math/amdgpu/fdim.cpp
new file mode 100644
index 00000000000000..f16942dc619341
--- /dev/null
+++ b/libc/src/math/amdgpu/fdim.cpp
@@ -0,0 +1,20 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/fdim.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, fdim, (double x, double y)) {
+ return __ocml_fdim_f64(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/fdimf.cpp b/libc/src/math/amdgpu/fdimf.cpp
new file mode 100644
index 00000000000000..eccb441f145561
--- /dev/null
+++ b/libc/src/math/amdgpu/fdimf.cpp
@@ -0,0 +1,20 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/fdimf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, fdimf, (float x, float y)) {
+ return __ocml_fdim_f32(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/floor.cpp b/libc/src/math/amdgpu/floor.cpp
similarity index 100%
rename from libc/src/math/gpu/floor.cpp
rename to libc/src/math/amdgpu/floor.cpp
diff --git a/libc/src/math/gpu/floorf.cpp b/libc/src/math/amdgpu/floorf.cpp
similarity index 100%
rename from libc/src/math/gpu/floorf.cpp
rename to libc/src/math/amdgpu/floorf.cpp
diff --git a/libc/src/math/gpu/fma.cpp b/libc/src/math/amdgpu/fma.cpp
similarity index 100%
rename from libc/src/math/gpu/fma.cpp
rename to libc/src/math/amdgpu/fma.cpp
diff --git a/libc/src/math/gpu/fmaf.cpp b/libc/src/math/amdgpu/fmaf.cpp
similarity index 100%
rename from libc/src/math/gpu/fmaf.cpp
rename to libc/src/math/amdgpu/fmaf.cpp
diff --git a/libc/src/math/gpu/fmax.cpp b/libc/src/math/amdgpu/fmax.cpp
similarity index 100%
rename from libc/src/math/gpu/fmax.cpp
rename to libc/src/math/amdgpu/fmax.cpp
diff --git a/libc/src/math/gpu/fmaxf.cpp b/libc/src/math/amdgpu/fmaxf.cpp
similarity index 100%
rename from libc/src/math/gpu/fmaxf.cpp
rename to libc/src/math/amdgpu/fmaxf.cpp
diff --git a/libc/src/math/gpu/fmin.cpp b/libc/src/math/amdgpu/fmin.cpp
similarity index 100%
rename from libc/src/math/gpu/fmin.cpp
rename to libc/src/math/amdgpu/fmin.cpp
diff --git a/libc/src/math/gpu/fminf.cpp b/libc/src/math/amdgpu/fminf.cpp
similarity index 100%
rename from libc/src/math/gpu/fminf.cpp
rename to libc/src/math/amdgpu/fminf.cpp
diff --git a/libc/src/math/gpu/fmod.cpp b/libc/src/math/amdgpu/fmod.cpp
similarity index 100%
rename from libc/src/math/gpu/fmod.cpp
rename to libc/src/math/amdgpu/fmod.cpp
diff --git a/libc/src/math/gpu/fmodf.cpp b/libc/src/math/amdgpu/fmodf.cpp
similarity index 100%
rename from libc/src/math/gpu/fmodf.cpp
rename to libc/src/math/amdgpu/fmodf.cpp
diff --git a/libc/src/math/amdgpu/frexp.cpp b/libc/src/math/amdgpu/frexp.cpp
new file mode 100644
index 00000000000000..0acf97342fa084
--- /dev/null
+++ b/libc/src/math/amdgpu/frexp.cpp
@@ -0,0 +1,20 @@
+//===-- 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 "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, frexp, (double x, int *p)) {
+ return __builtin_frexp(x, p);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/frexpf.cpp b/libc/src/math/amdgpu/frexpf.cpp
new file mode 100644
index 00000000000000..d870bf3095b25a
--- /dev/null
+++ b/libc/src/math/amdgpu/frexpf.cpp
@@ -0,0 +1,20 @@
+//===-- 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 "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, frexpf, (float x, int *p)) {
+ return __builtin_frexpf(x, p);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/hypot.cpp b/libc/src/math/amdgpu/hypot.cpp
new file mode 100644
index 00000000000000..ffc13504c8f418
--- /dev/null
+++ b/libc/src/math/amdgpu/hypot.cpp
@@ -0,0 +1,20 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/hypot.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, hypot, (double x, double y)) {
+ return __ocml_hypot_f64(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/hypotf.cpp b/libc/src/math/amdgpu/hypotf.cpp
new file mode 100644
index 00000000000000..811fc540488d3f
--- /dev/null
+++ b/libc/src/math/amdgpu/hypotf.cpp
@@ -0,0 +1,20 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/hypotf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, hypotf, (float x, float y)) {
+ return __ocml_hypot_f32(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/ilogb.cpp b/libc/src/math/amdgpu/ilogb.cpp
new file mode 100644
index 00000000000000..4479908d385648
--- /dev/null
+++ b/libc/src/math/amdgpu/ilogb.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/ilogb.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(int, ilogb, (double x)) { return __ocml_ilogb_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/ilogbf.cpp b/libc/src/math/amdgpu/ilogbf.cpp
new file mode 100644
index 00000000000000..cded285c72f38d
--- /dev/null
+++ b/libc/src/math/amdgpu/ilogbf.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/ilogbf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(int, ilogbf, (float x)) { return __ocml_ilogb_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/ldexp.cpp b/libc/src/math/amdgpu/ldexp.cpp
new file mode 100644
index 00000000000000..70c5b0d6e555ff
--- /dev/null
+++ b/libc/src/math/amdgpu/ldexp.cpp
@@ -0,0 +1,20 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/ldexp.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, ldexp, (double x, int y)) {
+ return __builtin_ldexp(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/ldexpf.cpp b/libc/src/math/amdgpu/ldexpf.cpp
new file mode 100644
index 00000000000000..8dc7c132ee213e
--- /dev/null
+++ b/libc/src/math/amdgpu/ldexpf.cpp
@@ -0,0 +1,20 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/ldexpf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, ldexpf, (float x, int y)) {
+ return __builtin_ldexpf(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/llrint.cpp b/libc/src/math/amdgpu/llrint.cpp
similarity index 87%
rename from libc/src/math/gpu/vendor/llrint.cpp
rename to libc/src/math/amdgpu/llrint.cpp
index aafd1609002b2b..307420a9b8b261 100644
--- a/libc/src/math/gpu/vendor/llrint.cpp
+++ b/libc/src/math/amdgpu/llrint.cpp
@@ -9,12 +9,12 @@
#include "src/math/llrint.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(long long, llrint, (double x)) {
- return internal::llrint(x);
+ return static_cast<long long>(__builtin_rint(x));
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/llrintf.cpp b/libc/src/math/amdgpu/llrintf.cpp
similarity index 87%
rename from libc/src/math/gpu/vendor/llrintf.cpp
rename to libc/src/math/amdgpu/llrintf.cpp
index 39cd3ad0c021a5..23404990fb1bdf 100644
--- a/libc/src/math/gpu/vendor/llrintf.cpp
+++ b/libc/src/math/amdgpu/llrintf.cpp
@@ -9,12 +9,12 @@
#include "src/math/llrintf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(long long, llrintf, (float x)) {
- return internal::llrintf(x);
+ return static_cast<long long>(__builtin_rintf(x));
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/llround.cpp b/libc/src/math/amdgpu/llround.cpp
similarity index 100%
rename from libc/src/math/gpu/llround.cpp
rename to libc/src/math/amdgpu/llround.cpp
diff --git a/libc/src/math/gpu/llroundf.cpp b/libc/src/math/amdgpu/llroundf.cpp
similarity index 100%
rename from libc/src/math/gpu/llroundf.cpp
rename to libc/src/math/amdgpu/llroundf.cpp
diff --git a/libc/src/math/amdgpu/log.cpp b/libc/src/math/amdgpu/log.cpp
new file mode 100644
index 00000000000000..3f2489580356f7
--- /dev/null
+++ b/libc/src/math/amdgpu/log.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU log function ----------------------------===//
+//
+// 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/log.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, log, (double x)) { return __ocml_log_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/log10.cpp b/libc/src/math/amdgpu/log10.cpp
new file mode 100644
index 00000000000000..d522d5e84291d2
--- /dev/null
+++ b/libc/src/math/amdgpu/log10.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU log10 function --------------------------===//
+//
+// 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/log10.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, log10, (double x)) { return __ocml_log10_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/log10f.cpp b/libc/src/math/amdgpu/log10f.cpp
new file mode 100644
index 00000000000000..47b9b162ad9db5
--- /dev/null
+++ b/libc/src/math/amdgpu/log10f.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU log10f function -------------------------===//
+//
+// 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/log10f.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, log10f, (float x)) { return __ocml_log10_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/log1p.cpp b/libc/src/math/amdgpu/log1p.cpp
new file mode 100644
index 00000000000000..fae60e44858611
--- /dev/null
+++ b/libc/src/math/amdgpu/log1p.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU log1p function --------------------------===//
+//
+// 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/log1p.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, log1p, (double x)) { return __ocml_log1p_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/log1pf.cpp b/libc/src/math/amdgpu/log1pf.cpp
new file mode 100644
index 00000000000000..e7b1772158fcda
--- /dev/null
+++ b/libc/src/math/amdgpu/log1pf.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU log1pf function -------------------------===//
+//
+// 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/log1pf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, log1pf, (float x)) { return __ocml_log1p_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/log2.cpp b/libc/src/math/amdgpu/log2.cpp
new file mode 100644
index 00000000000000..9d84f62dff6fcf
--- /dev/null
+++ b/libc/src/math/amdgpu/log2.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU log2 function ---------------------------===//
+//
+// 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/log2.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, log2, (double x)) { return __ocml_log2_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/log2f.cpp b/libc/src/math/amdgpu/log2f.cpp
new file mode 100644
index 00000000000000..7742a6141c37dc
--- /dev/null
+++ b/libc/src/math/amdgpu/log2f.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU log2f function --------------------------===//
+//
+// 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/log2f.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, log2f, (float x)) { return __ocml_log2_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/logb.cpp b/libc/src/math/amdgpu/logb.cpp
new file mode 100644
index 00000000000000..1344fbb182af60
--- /dev/null
+++ b/libc/src/math/amdgpu/logb.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU logb function ---------------------------===//
+//
+// 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/logb.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, logb, (double x)) { return __ocml_logb_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/logbf.cpp b/libc/src/math/amdgpu/logbf.cpp
new file mode 100644
index 00000000000000..fdb493fe28daf3
--- /dev/null
+++ b/libc/src/math/amdgpu/logbf.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU logbf function --------------------------===//
+//
+// 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/logbf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, logbf, (float x)) { return __ocml_logb_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/logf.cpp b/libc/src/math/amdgpu/logf.cpp
new file mode 100644
index 00000000000000..d4d4b265ee5cd7
--- /dev/null
+++ b/libc/src/math/amdgpu/logf.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU logf function ---------------------------===//
+//
+// 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/logf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, logf, (float x)) { return __ocml_log_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/lrint.cpp b/libc/src/math/amdgpu/lrint.cpp
new file mode 100644
index 00000000000000..b335b4f06393ca
--- /dev/null
+++ b/libc/src/math/amdgpu/lrint.cpp
@@ -0,0 +1,20 @@
+//===-- Implementation of the lrint 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/lrint.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(long, lrint, (double x)) {
+ return static_cast<long>(__builtin_rint(x));
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/lrintf.cpp b/libc/src/math/amdgpu/lrintf.cpp
new file mode 100644
index 00000000000000..7959e76728a434
--- /dev/null
+++ b/libc/src/math/amdgpu/lrintf.cpp
@@ -0,0 +1,20 @@
+//===-- Implementation of the lrintf 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/lrintf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(long, lrintf, (float x)) {
+ return static_cast<long>(__builtin_rintf(x));
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/lround.cpp b/libc/src/math/amdgpu/lround.cpp
similarity index 100%
rename from libc/src/math/gpu/lround.cpp
rename to libc/src/math/amdgpu/lround.cpp
diff --git a/libc/src/math/gpu/lroundf.cpp b/libc/src/math/amdgpu/lroundf.cpp
similarity index 100%
rename from libc/src/math/gpu/lroundf.cpp
rename to libc/src/math/amdgpu/lroundf.cpp
diff --git a/libc/src/math/gpu/modf.cpp b/libc/src/math/amdgpu/modf.cpp
similarity index 100%
rename from libc/src/math/gpu/modf.cpp
rename to libc/src/math/amdgpu/modf.cpp
diff --git a/libc/src/math/gpu/modff.cpp b/libc/src/math/amdgpu/modff.cpp
similarity index 100%
rename from libc/src/math/gpu/modff.cpp
rename to libc/src/math/amdgpu/modff.cpp
diff --git a/libc/src/math/gpu/nearbyint.cpp b/libc/src/math/amdgpu/nearbyint.cpp
similarity index 100%
rename from libc/src/math/gpu/nearbyint.cpp
rename to libc/src/math/amdgpu/nearbyint.cpp
diff --git a/libc/src/math/gpu/nearbyintf.cpp b/libc/src/math/amdgpu/nearbyintf.cpp
similarity index 100%
rename from libc/src/math/gpu/nearbyintf.cpp
rename to libc/src/math/amdgpu/nearbyintf.cpp
diff --git a/libc/src/math/amdgpu/nextafter.cpp b/libc/src/math/amdgpu/nextafter.cpp
new file mode 100644
index 00000000000000..5c74ef165268be
--- /dev/null
+++ b/libc/src/math/amdgpu/nextafter.cpp
@@ -0,0 +1,20 @@
+//===-- Implementation of the nextafter 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/nextafter.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, nextafter, (double x, double y)) {
+ return __ocml_nextafter_f64(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/nextafterf.cpp b/libc/src/math/amdgpu/nextafterf.cpp
new file mode 100644
index 00000000000000..a97b990a61fcd4
--- /dev/null
+++ b/libc/src/math/amdgpu/nextafterf.cpp
@@ -0,0 +1,20 @@
+//===-- Implementation of the nextafterf 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/nextafterf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, nextafterf, (float x, float y)) {
+ return __ocml_nextafter_f32(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/amdgpu/platform.h b/libc/src/math/amdgpu/platform.h
similarity index 100%
rename from libc/src/math/gpu/vendor/amdgpu/platform.h
rename to libc/src/math/amdgpu/platform.h
diff --git a/libc/src/math/gpu/vendor/pow.cpp b/libc/src/math/amdgpu/pow.cpp
similarity index 90%
rename from libc/src/math/gpu/vendor/pow.cpp
rename to libc/src/math/amdgpu/pow.cpp
index d49f2610a6919d..e5056f67292d53 100644
--- a/libc/src/math/gpu/vendor/pow.cpp
+++ b/libc/src/math/amdgpu/pow.cpp
@@ -9,12 +9,12 @@
#include "src/math/pow.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(double, pow, (double x, double y)) {
- return internal::pow(x, y);
+ return __ocml_pow_f64(x, y);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/powf.cpp b/libc/src/math/amdgpu/powf.cpp
similarity index 90%
rename from libc/src/math/gpu/vendor/powf.cpp
rename to libc/src/math/amdgpu/powf.cpp
index 37e02d252b74fd..6114bcc642e11b 100644
--- a/libc/src/math/gpu/vendor/powf.cpp
+++ b/libc/src/math/amdgpu/powf.cpp
@@ -9,12 +9,12 @@
#include "src/math/powf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(float, powf, (float x, float y)) {
- return internal::powf(x, y);
+ return __ocml_pow_f32(x, y);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/remainder.cpp b/libc/src/math/amdgpu/remainder.cpp
similarity index 100%
rename from libc/src/math/gpu/remainder.cpp
rename to libc/src/math/amdgpu/remainder.cpp
diff --git a/libc/src/math/gpu/remainderf.cpp b/libc/src/math/amdgpu/remainderf.cpp
similarity index 100%
rename from libc/src/math/gpu/remainderf.cpp
rename to libc/src/math/amdgpu/remainderf.cpp
diff --git a/libc/src/math/amdgpu/remquo.cpp b/libc/src/math/amdgpu/remquo.cpp
new file mode 100644
index 00000000000000..d8074a9626ecbb
--- /dev/null
+++ b/libc/src/math/amdgpu/remquo.cpp
@@ -0,0 +1,23 @@
+//===-- Implementation of the GPU remquo function -------------------------===//
+//
+// 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/remquo.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, remquo, (double x, double y, int *quo)) {
+ int tmp;
+ double r = __ocml_remquo_f64(x, y, (gpu::Private<int> *)&tmp);
+ *quo = tmp;
+ return r;
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/remquof.cpp b/libc/src/math/amdgpu/remquof.cpp
new file mode 100644
index 00000000000000..b6584dfb97c3b5
--- /dev/null
+++ b/libc/src/math/amdgpu/remquof.cpp
@@ -0,0 +1,23 @@
+//===-- Implementation of the GPU remquof function ------------------------===//
+//
+// 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/remquof.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, remquof, (float x, float y, int *quo)) {
+ int tmp;
+ float r = __ocml_remquo_f32(x, y, (gpu::Private<int> *)&tmp);
+ *quo = tmp;
+ return r;
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/rint.cpp b/libc/src/math/amdgpu/rint.cpp
similarity index 100%
rename from libc/src/math/gpu/rint.cpp
rename to libc/src/math/amdgpu/rint.cpp
diff --git a/libc/src/math/gpu/rintf.cpp b/libc/src/math/amdgpu/rintf.cpp
similarity index 100%
rename from libc/src/math/gpu/rintf.cpp
rename to libc/src/math/amdgpu/rintf.cpp
diff --git a/libc/src/math/gpu/round.cpp b/libc/src/math/amdgpu/round.cpp
similarity index 100%
rename from libc/src/math/gpu/round.cpp
rename to libc/src/math/amdgpu/round.cpp
diff --git a/libc/src/math/gpu/roundf.cpp b/libc/src/math/amdgpu/roundf.cpp
similarity index 100%
rename from libc/src/math/gpu/roundf.cpp
rename to libc/src/math/amdgpu/roundf.cpp
diff --git a/libc/src/math/amdgpu/scalbn.cpp b/libc/src/math/amdgpu/scalbn.cpp
new file mode 100644
index 00000000000000..c2a43e03a7bcd8
--- /dev/null
+++ b/libc/src/math/amdgpu/scalbn.cpp
@@ -0,0 +1,20 @@
+//===-- Implementation of the GPU scalbn function -------------------------===//
+//
+// 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/scalbn.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, scalbn, (double x, int y)) {
+ return __builtin_amdgcn_ldexp(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/scalbnf.cpp b/libc/src/math/amdgpu/scalbnf.cpp
new file mode 100644
index 00000000000000..63de26ccbc4704
--- /dev/null
+++ b/libc/src/math/amdgpu/scalbnf.cpp
@@ -0,0 +1,20 @@
+//===-- Implementation of the GPU scalbnf function ------------------------===//
+//
+// 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/scalbnf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, scalbnf, (float x, int y)) {
+ return __builtin_amdgcn_ldexpf(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/sin.cpp b/libc/src/math/amdgpu/sin.cpp
new file mode 100644
index 00000000000000..dbc29a725db076
--- /dev/null
+++ b/libc/src/math/amdgpu/sin.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the sin 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/sin.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, sin, (double x)) { return __ocml_sin_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/sincos.cpp b/libc/src/math/amdgpu/sincos.cpp
new file mode 100644
index 00000000000000..7cdd0d1f97690a
--- /dev/null
+++ b/libc/src/math/amdgpu/sincos.cpp
@@ -0,0 +1,20 @@
+//===-- Implementation of the sincos 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/sincos.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(void, sincos, (double x, double *sinptr, double *cosptr)) {
+ *sinptr = __ocml_sincos_f64(x, cosptr);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/sincosf.cpp b/libc/src/math/amdgpu/sincosf.cpp
similarity index 89%
rename from libc/src/math/gpu/vendor/sincosf.cpp
rename to libc/src/math/amdgpu/sincosf.cpp
index 17892ceb107da5..37a5e2a6d11c2d 100644
--- a/libc/src/math/gpu/vendor/sincosf.cpp
+++ b/libc/src/math/amdgpu/sincosf.cpp
@@ -9,12 +9,12 @@
#include "src/math/sincosf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(void, sincosf, (float x, float *sinptr, float *cosptr)) {
- return internal::sincosf(x, sinptr, cosptr);
+ *sinptr = __ocml_sincos_f32(x, cosptr);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/sinf.cpp b/libc/src/math/amdgpu/sinf.cpp
new file mode 100644
index 00000000000000..cda2c626f8ac67
--- /dev/null
+++ b/libc/src/math/amdgpu/sinf.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the sinf 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/sinf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, sinf, (float x)) { return __ocml_sin_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/sinh.cpp b/libc/src/math/amdgpu/sinh.cpp
similarity index 71%
rename from libc/src/math/gpu/sinh.cpp
rename to libc/src/math/amdgpu/sinh.cpp
index 054e046f2abd3b..66cacd19e4632a 100644
--- a/libc/src/math/gpu/sinh.cpp
+++ b/libc/src/math/amdgpu/sinh.cpp
@@ -1,4 +1,4 @@
-//===-- Implementation of the GPU sinh function ---------------------------===//
+//===-- Implementation of the sinh 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.
@@ -9,8 +9,10 @@
#include "src/math/sinh.h"
#include "src/__support/common.h"
+#include "declarations.h"
+
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, sinh, (double x)) { return __builtin_sinh(x); }
+LLVM_LIBC_FUNCTION(double, sinh, (double x)) { return __ocml_sinh_f64(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/sinhf.cpp b/libc/src/math/amdgpu/sinhf.cpp
new file mode 100644
index 00000000000000..5d3f5ea6b36ae4
--- /dev/null
+++ b/libc/src/math/amdgpu/sinhf.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the sinhf 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/sinhf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, sinhf, (float x)) { return __ocml_sinh_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/sqrt.cpp b/libc/src/math/amdgpu/sqrt.cpp
similarity index 100%
rename from libc/src/math/gpu/sqrt.cpp
rename to libc/src/math/amdgpu/sqrt.cpp
diff --git a/libc/src/math/gpu/sqrtf.cpp b/libc/src/math/amdgpu/sqrtf.cpp
similarity index 100%
rename from libc/src/math/gpu/sqrtf.cpp
rename to libc/src/math/amdgpu/sqrtf.cpp
diff --git a/libc/src/math/gpu/tan.cpp b/libc/src/math/amdgpu/tan.cpp
similarity index 72%
rename from libc/src/math/gpu/tan.cpp
rename to libc/src/math/amdgpu/tan.cpp
index d02b1063561246..6121a9319a2abd 100644
--- a/libc/src/math/gpu/tan.cpp
+++ b/libc/src/math/amdgpu/tan.cpp
@@ -1,4 +1,4 @@
-//===-- Implementation of the GPU tan function ----------------------------===//
+//===-- Implementation of the tan 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.
@@ -9,8 +9,10 @@
#include "src/math/tan.h"
#include "src/__support/common.h"
+#include "declarations.h"
+
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, tan, (double x)) { return __builtin_tan(x); }
+LLVM_LIBC_FUNCTION(double, tan, (double x)) { return __ocml_tan_f64(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/tanf.cpp b/libc/src/math/amdgpu/tanf.cpp
new file mode 100644
index 00000000000000..fdd83ee7aeb942
--- /dev/null
+++ b/libc/src/math/amdgpu/tanf.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the tanf 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/tanf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, tanf, (float x)) { return __ocml_tan_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/tanh.cpp b/libc/src/math/amdgpu/tanh.cpp
similarity index 71%
rename from libc/src/math/gpu/tanh.cpp
rename to libc/src/math/amdgpu/tanh.cpp
index 778e883acba044..25a9c2954bf39c 100644
--- a/libc/src/math/gpu/tanh.cpp
+++ b/libc/src/math/amdgpu/tanh.cpp
@@ -1,4 +1,4 @@
-//===-- Implementation of the GPU tanh function ---------------------------===//
+//===-- Implementation of the tanh 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.
@@ -9,8 +9,10 @@
#include "src/math/tanh.h"
#include "src/__support/common.h"
+#include "declarations.h"
+
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, tanh, (double x)) { return __builtin_tanh(x); }
+LLVM_LIBC_FUNCTION(double, tanh, (double x)) { return __ocml_tanh_f64(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/tanhf.cpp b/libc/src/math/amdgpu/tanhf.cpp
new file mode 100644
index 00000000000000..a4bfd2095eb352
--- /dev/null
+++ b/libc/src/math/amdgpu/tanhf.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the tanhf 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/tanhf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, tanhf, (float x)) { return __ocml_tanh_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/tgamma.cpp b/libc/src/math/amdgpu/tgamma.cpp
new file mode 100644
index 00000000000000..10f58d566f88aa
--- /dev/null
+++ b/libc/src/math/amdgpu/tgamma.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU tgamma function -------------------------===//
+//
+// 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/tgamma.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, tgamma, (double x)) { return __ocml_tgamma_f64(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/amdgpu/tgammaf.cpp b/libc/src/math/amdgpu/tgammaf.cpp
new file mode 100644
index 00000000000000..e7d22059a7c41e
--- /dev/null
+++ b/libc/src/math/amdgpu/tgammaf.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU tgammaf function ------------------------===//
+//
+// 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/tgammaf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, tgammaf, (float x)) { return __ocml_tgamma_f32(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/trunc.cpp b/libc/src/math/amdgpu/trunc.cpp
similarity index 100%
rename from libc/src/math/gpu/trunc.cpp
rename to libc/src/math/amdgpu/trunc.cpp
diff --git a/libc/src/math/gpu/truncf.cpp b/libc/src/math/amdgpu/truncf.cpp
similarity index 100%
rename from libc/src/math/gpu/truncf.cpp
rename to libc/src/math/amdgpu/truncf.cpp
diff --git a/libc/src/math/gpu/CMakeLists.txt b/libc/src/math/gpu/CMakeLists.txt
deleted file mode 100644
index 75a916e2a0112a..00000000000000
--- a/libc/src/math/gpu/CMakeLists.txt
+++ /dev/null
@@ -1,384 +0,0 @@
-# Math functions not yet available in the libc project, or those not yet tuned
-# for GPU workloads are provided as wrappers over vendor libraries. If we find
-# them ahead of time we will import them statically. Otherwise, we will keep
-# them as external references and expect them to be resolved by the user when
-# they compile. In the future,we will use implementations from the 'libc'
-# project and not provide these wrappers.
-add_subdirectory(vendor)
-
-# For the GPU we want to be able to optionally depend on the vendor libraries
-# until we have a suitable replacement inside `libc`.
-# TODO: We should have an option to enable or disable these on a per-function
-# basis.
-option(LIBC_GPU_VENDOR_MATH "Use vendor wrappers for GPU math" ON)
-function(add_math_entrypoint_gpu_object name)
- get_fq_target_name("vendor.${name}" fq_vendor_specific_target_name)
- if(TARGET ${fq_vendor_specific_target_name} AND ${LIBC_GPU_VENDOR_MATH})
- return()
- endif()
-
- 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(
- 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(
- 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(
- 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(
- 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(
- 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(
- 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(
- lround
- SRCS
- lround.cpp
- HDRS
- ../lround.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- lroundf
- SRCS
- lroundf.cpp
- HDRS
- ../lroundf.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- llround
- SRCS
- llround.cpp
- HDRS
- ../llround.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- llroundf
- SRCS
- llroundf.cpp
- HDRS
- ../llroundf.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- modf
- SRCS
- modf.cpp
- HDRS
- ../modf.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- modff
- SRCS
- modff.cpp
- HDRS
- ../modff.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- nearbyint
- SRCS
- nearbyint.cpp
- HDRS
- ../nearbyint.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- nearbyintf
- SRCS
- nearbyintf.cpp
- HDRS
- ../nearbyintf.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- remainder
- SRCS
- remainder.cpp
- HDRS
- ../remainder.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- remainderf
- SRCS
- remainderf.cpp
- HDRS
- ../remainderf.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- rint
- SRCS
- rint.cpp
- HDRS
- ../rint.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- rintf
- SRCS
- rintf.cpp
- HDRS
- ../rintf.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- round
- SRCS
- round.cpp
- HDRS
- ../round.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- sinh
- SRCS
- sinh.cpp
- HDRS
- ../sinh.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- sqrt
- SRCS
- sqrt.cpp
- HDRS
- ../sqrt.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- sqrtf
- SRCS
- sqrtf.cpp
- HDRS
- ../sqrtf.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- tan
- SRCS
- tan.cpp
- HDRS
- ../tan.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- tanh
- SRCS
- tanh.cpp
- HDRS
- ../tanh.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- trunc
- SRCS
- trunc.cpp
- HDRS
- ../trunc.h
- COMPILE_OPTIONS
- -O2
-)
-
-add_math_entrypoint_gpu_object(
- truncf
- SRCS
- truncf.cpp
- HDRS
- ../truncf.h
- COMPILE_OPTIONS
- -O2
-)
diff --git a/libc/src/math/gpu/vendor/amdgpu/amdgpu.h b/libc/src/math/gpu/vendor/amdgpu/amdgpu.h
deleted file mode 100644
index 43961fc75982ab..00000000000000
--- a/libc/src/math/gpu/vendor/amdgpu/amdgpu.h
+++ /dev/null
@@ -1,127 +0,0 @@
-//===-- AMDGPU specific definitions 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_H
-#define LLVM_LIBC_SRC_MATH_GPU_AMDGPU_H
-
-#include "declarations.h"
-#include "platform.h"
-
-#include "src/__support/macros/attributes.h"
-
-namespace LIBC_NAMESPACE {
-namespace internal {
-LIBC_INLINE double acos(double x) { return __ocml_acos_f64(x); }
-LIBC_INLINE float acosf(float x) { return __ocml_acos_f32(x); }
-LIBC_INLINE double acosh(double x) { return __ocml_acosh_f64(x); }
-LIBC_INLINE float acoshf(float x) { return __ocml_acosh_f32(x); }
-LIBC_INLINE double asin(double x) { return __ocml_asin_f64(x); }
-LIBC_INLINE float asinf(float x) { return __ocml_asin_f32(x); }
-LIBC_INLINE double asinh(double x) { return __ocml_asinh_f64(x); }
-LIBC_INLINE float asinhf(float x) { return __ocml_asinh_f32(x); }
-LIBC_INLINE double atan(double x) { return __ocml_atan_f64(x); }
-LIBC_INLINE float atanf(float x) { return __ocml_atan_f32(x); }
-LIBC_INLINE double atan2(double x, double y) { return __ocml_atan2_f64(x, y); }
-LIBC_INLINE float atan2f(float x, float y) { return __ocml_atan2_f32(x, y); }
-LIBC_INLINE double atanh(double x) { return __ocml_atanh_f64(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 double cosh(double x) { return __ocml_cosh_f64(x); }
-LIBC_INLINE float coshf(float x) { return __ocml_cosh_f32(x); }
-LIBC_INLINE double erf(double x) { return __ocml_erf_f64(x); }
-LIBC_INLINE float erff(float x) { return __ocml_erf_f32(x); }
-LIBC_INLINE double exp(double x) { return __builtin_exp(x); }
-LIBC_INLINE float expf(float x) { return __builtin_expf(x); }
-LIBC_INLINE double exp2(double x) { return __ocml_exp2_f64(x); }
-LIBC_INLINE float exp2f(float x) { return __ocml_exp2_f32(x); }
-LIBC_INLINE double exp10(double x) { return __ocml_exp10_f64(x); }
-LIBC_INLINE float exp10f(float x) { return __ocml_exp10_f32(x); }
-LIBC_INLINE double expm1(double x) { return __ocml_expm1_f64(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 __builtin_ldexp(x, i); }
-LIBC_INLINE float ldexpf(float x, int i) { return __builtin_ldexpf(x, i); }
-LIBC_INLINE long long llrint(double x) {
- return static_cast<long long>(__builtin_rint(x));
-}
-LIBC_INLINE long long llrintf(float x) {
- return static_cast<long long>(__builtin_rintf(x));
-}
-LIBC_INLINE double log10(double x) { return __ocml_log10_f64(x); }
-LIBC_INLINE float log10f(float x) { return __ocml_log10_f32(x); }
-LIBC_INLINE double log1p(double x) { return __ocml_log1p_f64(x); }
-LIBC_INLINE float log1pf(float x) { return __ocml_log1p_f32(x); }
-LIBC_INLINE double log2(double x) { return __ocml_log2_f64(x); }
-LIBC_INLINE float log2f(float x) { return __ocml_log2_f32(x); }
-LIBC_INLINE double log(double x) { return __ocml_log_f64(x); }
-LIBC_INLINE float logf(float x) { return __ocml_log_f32(x); }
-LIBC_INLINE long lrint(double x) {
- return static_cast<long>(__builtin_rint(x));
-}
-LIBC_INLINE long lrintf(float x) {
- return static_cast<long>(__builtin_rintf(x));
-}
-LIBC_INLINE double nextafter(double x, double y) {
- return __ocml_nextafter_f64(x, y);
-}
-LIBC_INLINE float nextafterf(float x, float y) {
- return __ocml_nextafter_f32(x, y);
-}
-LIBC_INLINE double pow(double x, double y) { return __ocml_pow_f64(x, y); }
-LIBC_INLINE float powf(float x, float y) { return __ocml_pow_f32(x, y); }
-LIBC_INLINE double sin(double x) { return __ocml_sin_f64(x); }
-LIBC_INLINE float sinf(float x) { return __ocml_sin_f32(x); }
-LIBC_INLINE void sincos(double x, double *sinptr, double *cosptr) {
- *sinptr = __ocml_sincos_f64(x, cosptr);
-}
-LIBC_INLINE void sincosf(float x, float *sinptr, float *cosptr) {
- *sinptr = __ocml_sincos_f32(x, cosptr);
-}
-LIBC_INLINE double sinh(double x) { return __ocml_sinh_f64(x); }
-LIBC_INLINE float sinhf(float x) { return __ocml_sinh_f32(x); }
-LIBC_INLINE double tan(double x) { return __ocml_tan_f64(x); }
-LIBC_INLINE float tanf(float x) { return __ocml_tan_f32(x); }
-LIBC_INLINE double tanh(double x) { return __ocml_tanh_f64(x); }
-LIBC_INLINE float tanhf(float x) { return __ocml_tanh_f32(x); }
-LIBC_INLINE double scalbn(double x, int i) {
- return __builtin_amdgcn_ldexp(x, i);
-}
-LIBC_INLINE float scalbnf(float x, int i) {
- return __builtin_amdgcn_ldexpf(x, i);
-}
-LIBC_INLINE double frexp(double x, int *nptr) {
- return __builtin_frexp(x, nptr);
-}
-LIBC_INLINE float frexpf(float x, int *nptr) {
- return __builtin_frexpf(x, nptr);
-}
-LIBC_INLINE double remquo(double x, double y, int *q) {
- int tmp;
- double r = __ocml_remquo_f64(x, y, (gpu::Private<int> *)&tmp);
- *q = tmp;
- return r;
-}
-LIBC_INLINE float remquof(float x, float y, int *q) {
- int tmp;
- float r = __ocml_remquo_f32(x, y, (gpu::Private<int> *)&tmp);
- *q = tmp;
- return r;
-}
-LIBC_INLINE double tgamma(double x) { return __ocml_tgamma_f64(x); }
-LIBC_INLINE float tgammaf(float x) { return __ocml_tgamma_f32(x); }
-
-} // namespace internal
-} // namespace LIBC_NAMESPACE
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_AMDGPU_H
diff --git a/libc/src/math/gpu/vendor/common.h b/libc/src/math/gpu/vendor/common.h
deleted file mode 100644
index 041a9a01c30e93..00000000000000
--- a/libc/src/math/gpu/vendor/common.h
+++ /dev/null
@@ -1,22 +0,0 @@
-//===-- Common interface for compiling the GPU math -----------------------===//
-//
-// 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_COMMON_H
-#define LLVM_LIBC_SRC_MATH_GPU_COMMON_H
-
-#include "src/__support/macros/properties/architectures.h"
-
-#if defined(LIBC_TARGET_ARCH_IS_AMDGPU)
-#include "amdgpu/amdgpu.h"
-#elif defined(LIBC_TARGET_ARCH_IS_NVPTX)
-#include "nvptx/nvptx.h"
-#else
-#error "Unsupported platform"
-#endif
-
-#endif // LLVM_LIBC_SRC_MATH_GPU_COMMON_H
diff --git a/libc/src/math/nvptx/CMakeLists.txt b/libc/src/math/nvptx/CMakeLists.txt
new file mode 100644
index 00000000000000..194e1fa7af491e
--- /dev/null
+++ b/libc/src/math/nvptx/CMakeLists.txt
@@ -0,0 +1,1179 @@
+# Math functions not yet available in the libc project, or those not yet tuned
+# for GPU workloads are provided as wrappers over vendor libraries. If we find
+# them ahead of time we will import them statically. Otherwise, we will keep
+# them as external references and expect them to be resolved by the user when
+# they compile. In the future,we will use implementations from the 'libc'
+# project and not provide these wrappers.
+if(CUDAToolkit_FOUND)
+ set(libdevice_path ${CUDAToolkit_BIN_DIR}/../nvvm/libdevice/libdevice.10.bc)
+ if (EXISTS ${libdevice_path})
+ message(STATUS "Found the CUDA device library. Implementations falling back "
+ "to the vendor libraries will be resolved statically.")
+ set(bitcode_link_flags
+ "SHELL:-Xclang -mlink-builtin-bitcode -Xclang ${libdevice_path}")
+ endif()
+else()
+ message(STATUS "Could not find the CUDA device library. Unimplemented "
+ "functions will be an external reference to the vendor libraries.")
+endif()
+
+add_entrypoint_object(
+ ceil
+ SRCS
+ ceil.cpp
+ HDRS
+ ../ceil.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ ceilf
+ SRCS
+ ceilf.cpp
+ HDRS
+ ../ceilf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ copysign
+ SRCS
+ copysign.cpp
+ HDRS
+ ../copysign.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ copysignf
+ SRCS
+ copysignf.cpp
+ HDRS
+ ../copysignf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fabs
+ SRCS
+ fabs.cpp
+ HDRS
+ ../fabs.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fabsf
+ SRCS
+ fabsf.cpp
+ HDRS
+ ../fabsf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ floor
+ SRCS
+ floor.cpp
+ HDRS
+ ../floor.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ floorf
+ SRCS
+ floorf.cpp
+ HDRS
+ ../floorf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fma
+ SRCS
+ fma.cpp
+ HDRS
+ ../fma.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fmaf
+ SRCS
+ fmaf.cpp
+ HDRS
+ ../fmaf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fmax
+ SRCS
+ fmax.cpp
+ HDRS
+ ../fmax.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fmaxf
+ SRCS
+ fmaxf.cpp
+ HDRS
+ ../fmaxf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fmin
+ SRCS
+ fmin.cpp
+ HDRS
+ ../fmin.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fminf
+ SRCS
+ fminf.cpp
+ HDRS
+ ../fminf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fmod
+ SRCS
+ fmod.cpp
+ HDRS
+ ../fmod.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ fmodf
+ SRCS
+ fmodf.cpp
+ HDRS
+ ../fmodf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ lround
+ SRCS
+ lround.cpp
+ HDRS
+ ../lround.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ lroundf
+ SRCS
+ lroundf.cpp
+ HDRS
+ ../lroundf.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(
+ modf
+ SRCS
+ modf.cpp
+ HDRS
+ ../modf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ modff
+ SRCS
+ modff.cpp
+ HDRS
+ ../modff.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ nearbyint
+ SRCS
+ nearbyint.cpp
+ HDRS
+ ../nearbyint.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ nearbyintf
+ SRCS
+ nearbyintf.cpp
+ HDRS
+ ../nearbyintf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ remainder
+ SRCS
+ remainder.cpp
+ HDRS
+ ../remainder.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ remainderf
+ SRCS
+ remainderf.cpp
+ HDRS
+ ../remainderf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ rint
+ SRCS
+ rint.cpp
+ HDRS
+ ../rint.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ rintf
+ SRCS
+ rintf.cpp
+ HDRS
+ ../rintf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ round
+ SRCS
+ round.cpp
+ HDRS
+ ../round.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ sqrt
+ SRCS
+ sqrt.cpp
+ HDRS
+ ../sqrt.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ sqrtf
+ SRCS
+ sqrtf.cpp
+ HDRS
+ ../sqrtf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ trunc
+ SRCS
+ trunc.cpp
+ HDRS
+ ../trunc.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ truncf
+ SRCS
+ truncf.cpp
+ HDRS
+ ../truncf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+# The following functions currently are not implemented natively and borrow from
+# existing implementations. This will be removed in the future.
+add_entrypoint_object(
+ acos
+ SRCS
+ acos.cpp
+ HDRS
+ ../acos.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ acosf
+ SRCS
+ acosf.cpp
+ HDRS
+ ../acosf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ acosh
+ SRCS
+ acosh.cpp
+ HDRS
+ ../acosh.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ acoshf
+ SRCS
+ acoshf.cpp
+ HDRS
+ ../acoshf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ asin
+ SRCS
+ asin.cpp
+ HDRS
+ ../asin.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ asinf
+ SRCS
+ asinf.cpp
+ HDRS
+ ../asinf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ asinh
+ SRCS
+ asinh.cpp
+ HDRS
+ ../asinh.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ atan
+ SRCS
+ atan.cpp
+ HDRS
+ ../atan.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ atanf
+ SRCS
+ atanf.cpp
+ HDRS
+ ../atanf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ atan2
+ SRCS
+ atan2.cpp
+ HDRS
+ ../atan2.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ atan2f
+ SRCS
+ atan2f.cpp
+ HDRS
+ ../atan2f.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ atanh
+ SRCS
+ atanh.cpp
+ HDRS
+ ../atanh.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ atanhf
+ SRCS
+ atanhf.cpp
+ HDRS
+ ../atanhf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ cos
+ SRCS
+ cos.cpp
+ HDRS
+ ../cos.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ cosf
+ SRCS
+ cosf.cpp
+ HDRS
+ ../cosf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ cosh
+ SRCS
+ cosh.cpp
+ HDRS
+ ../cosh.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ coshf
+ SRCS
+ coshf.cpp
+ HDRS
+ ../coshf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ erf
+ SRCS
+ erf.cpp
+ HDRS
+ ../erf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ erff
+ SRCS
+ erff.cpp
+ HDRS
+ ../erff.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ exp
+ SRCS
+ exp.cpp
+ HDRS
+ ../exp.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ exp10
+ SRCS
+ exp10.cpp
+ HDRS
+ ../exp10.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ exp10f
+ SRCS
+ exp10f.cpp
+ HDRS
+ ../exp10f.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ exp2
+ SRCS
+ exp2.cpp
+ HDRS
+ ../exp2.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ exp2f
+ SRCS
+ exp2f.cpp
+ HDRS
+ ../exp2f.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ expf
+ SRCS
+ expf.cpp
+ HDRS
+ ../expf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ expm1
+ SRCS
+ expm1.cpp
+ HDRS
+ ../expm1.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ expm1f
+ SRCS
+ expm1f.cpp
+ HDRS
+ ../expm1f.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ fdim
+ SRCS
+ fdim.cpp
+ HDRS
+ ../fdim.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ fdimf
+ SRCS
+ fdimf.cpp
+ HDRS
+ ../fdimf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ hypot
+ SRCS
+ hypot.cpp
+ HDRS
+ ../hypot.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ hypotf
+ SRCS
+ hypotf.cpp
+ HDRS
+ ../hypotf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ ilogb
+ SRCS
+ ilogb.cpp
+ HDRS
+ ../ilogb.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ ilogbf
+ SRCS
+ ilogbf.cpp
+ HDRS
+ ../ilogbf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ log10
+ SRCS
+ log10.cpp
+ HDRS
+ ../log10.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ log10f
+ SRCS
+ log10f.cpp
+ HDRS
+ ../log10f.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ log2
+ SRCS
+ log2.cpp
+ HDRS
+ ../log2.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ log2f
+ SRCS
+ log2f.cpp
+ HDRS
+ ../log2f.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ log
+ SRCS
+ log.cpp
+ HDRS
+ ../log.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ logf
+ SRCS
+ logf.cpp
+ HDRS
+ ../logf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ lrint
+ SRCS
+ lrint.cpp
+ HDRS
+ ../lrint.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ lrintf
+ SRCS
+ lrintf.cpp
+ HDRS
+ ../lrintf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ ldexp
+ SRCS
+ ldexp.cpp
+ HDRS
+ ../ldexp.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ ldexpf
+ SRCS
+ ldexpf.cpp
+ HDRS
+ ../ldexpf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ log1p
+ SRCS
+ log1p.cpp
+ HDRS
+ ../log1p.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ log1pf
+ SRCS
+ log1pf.cpp
+ HDRS
+ ../log1pf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ llrint
+ SRCS
+ llrint.cpp
+ HDRS
+ ../llrint.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ llrintf
+ SRCS
+ llrintf.cpp
+ HDRS
+ ../llrintf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ remquo
+ SRCS
+ remquo.cpp
+ HDRS
+ ../remquo.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ remquof
+ SRCS
+ remquof.cpp
+ HDRS
+ ../remquof.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ scalbn
+ SRCS
+ scalbn.cpp
+ HDRS
+ ../scalbn.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ scalbnf
+ SRCS
+ scalbnf.cpp
+ HDRS
+ ../scalbnf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+
+add_entrypoint_object(
+ nextafter
+ SRCS
+ nextafter.cpp
+ HDRS
+ ../nextafter.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ nextafterf
+ SRCS
+ nextafterf.cpp
+ HDRS
+ ../nextafterf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ pow
+ SRCS
+ pow.cpp
+ HDRS
+ ../pow.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ powf
+ SRCS
+ powf.cpp
+ HDRS
+ ../powf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ sin
+ SRCS
+ sin.cpp
+ HDRS
+ ../sin.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ sinf
+ SRCS
+ sinf.cpp
+ HDRS
+ ../sinf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ sincos
+ SRCS
+ sincos.cpp
+ HDRS
+ ../sincos.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ sincosf
+ SRCS
+ sincosf.cpp
+ HDRS
+ ../sincosf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ sinh
+ SRCS
+ sinh.cpp
+ HDRS
+ ../sinh.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ sinhf
+ SRCS
+ sinhf.cpp
+ HDRS
+ ../sinhf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ tan
+ SRCS
+ tan.cpp
+ HDRS
+ ../tan.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ tanf
+ SRCS
+ tanf.cpp
+ HDRS
+ ../tanf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ tanh
+ SRCS
+ tanh.cpp
+ HDRS
+ ../tanh.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ tanhf
+ SRCS
+ tanhf.cpp
+ HDRS
+ ../tanhf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ tgamma
+ SRCS
+ tgamma.cpp
+ HDRS
+ ../tgamma.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ tgammaf
+ SRCS
+ tgammaf.cpp
+ HDRS
+ ../tgammaf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ frexp
+ SRCS
+ frexp.cpp
+ HDRS
+ ../frexp.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ frexpf
+ SRCS
+ frexpf.cpp
+ HDRS
+ ../frexpf.h
+ COMPILE_OPTIONS
+ ${itcode_link_flags}
+ -O2
+ VENDOR
+)
diff --git a/libc/src/math/gpu/vendor/acos.cpp b/libc/src/math/nvptx/acos.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/acos.cpp
rename to libc/src/math/nvptx/acos.cpp
index 83b674fa6ae3b7..da2c7952feba56 100644
--- a/libc/src/math/gpu/vendor/acos.cpp
+++ b/libc/src/math/nvptx/acos.cpp
@@ -9,10 +9,10 @@
#include "src/math/acos.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, acos, (double x)) { return internal::acos(x); }
+LLVM_LIBC_FUNCTION(double, acos, (double x)) { return __nv_acos(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/acosf.cpp b/libc/src/math/nvptx/acosf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/acosf.cpp
rename to libc/src/math/nvptx/acosf.cpp
index ac629761a4398a..8a4125f03b8cd7 100644
--- a/libc/src/math/gpu/vendor/acosf.cpp
+++ b/libc/src/math/nvptx/acosf.cpp
@@ -9,10 +9,10 @@
#include "src/math/acosf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, acosf, (float x)) { return internal::acosf(x); }
+LLVM_LIBC_FUNCTION(float, acosf, (float x)) { return __nv_acosf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/acosh.cpp b/libc/src/math/nvptx/acosh.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/acosh.cpp
rename to libc/src/math/nvptx/acosh.cpp
index cc1b8b572b3db7..06f6e2922f5693 100644
--- a/libc/src/math/gpu/vendor/acosh.cpp
+++ b/libc/src/math/nvptx/acosh.cpp
@@ -9,10 +9,10 @@
#include "src/math/acosh.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, acosh, (double x)) { return internal::acosh(x); }
+LLVM_LIBC_FUNCTION(double, acosh, (double x)) { return __nv_acosh(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/acoshf.cpp b/libc/src/math/nvptx/acoshf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/acoshf.cpp
rename to libc/src/math/nvptx/acoshf.cpp
index a0384f89eed3a6..00e8053a507836 100644
--- a/libc/src/math/gpu/vendor/acoshf.cpp
+++ b/libc/src/math/nvptx/acoshf.cpp
@@ -9,10 +9,10 @@
#include "src/math/acoshf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, acoshf, (float x)) { return internal::acoshf(x); }
+LLVM_LIBC_FUNCTION(float, acoshf, (float x)) { return __nv_acoshf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/asin.cpp b/libc/src/math/nvptx/asin.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/asin.cpp
rename to libc/src/math/nvptx/asin.cpp
index 24a8a136e88eb2..74d92fded72b96 100644
--- a/libc/src/math/gpu/vendor/asin.cpp
+++ b/libc/src/math/nvptx/asin.cpp
@@ -9,10 +9,10 @@
#include "src/math/asin.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, asin, (double x)) { return internal::asin(x); }
+LLVM_LIBC_FUNCTION(double, asin, (double x)) { return __nv_asin(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/asinf.cpp b/libc/src/math/nvptx/asinf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/asinf.cpp
rename to libc/src/math/nvptx/asinf.cpp
index 595a48f82744d5..30544bc1313d60 100644
--- a/libc/src/math/gpu/vendor/asinf.cpp
+++ b/libc/src/math/nvptx/asinf.cpp
@@ -9,10 +9,10 @@
#include "src/math/asinf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, asinf, (float x)) { return internal::asinf(x); }
+LLVM_LIBC_FUNCTION(float, asinf, (float x)) { return __nv_asinf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/asinh.cpp b/libc/src/math/nvptx/asinh.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/asinh.cpp
rename to libc/src/math/nvptx/asinh.cpp
index f417d9fd8c1f73..0e5dbb47e667e5 100644
--- a/libc/src/math/gpu/vendor/asinh.cpp
+++ b/libc/src/math/nvptx/asinh.cpp
@@ -9,10 +9,10 @@
#include "src/math/asinh.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, asinh, (double x)) { return internal::asinh(x); }
+LLVM_LIBC_FUNCTION(double, asinh, (double x)) { return __nv_asinh(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/asinhf.cpp b/libc/src/math/nvptx/asinhf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/asinhf.cpp
rename to libc/src/math/nvptx/asinhf.cpp
index 78e5543cf36656..6648108646cd36 100644
--- a/libc/src/math/gpu/vendor/asinhf.cpp
+++ b/libc/src/math/nvptx/asinhf.cpp
@@ -9,10 +9,10 @@
#include "src/math/asinhf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, asinhf, (float x)) { return internal::asinhf(x); }
+LLVM_LIBC_FUNCTION(float, asinhf, (float x)) { return __nv_asinhf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/atan.cpp b/libc/src/math/nvptx/atan.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/atan.cpp
rename to libc/src/math/nvptx/atan.cpp
index 45d7f02c02cd63..3af793a53ae5e4 100644
--- a/libc/src/math/gpu/vendor/atan.cpp
+++ b/libc/src/math/nvptx/atan.cpp
@@ -9,10 +9,10 @@
#include "src/math/atan.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, atan, (double x)) { return internal::atan(x); }
+LLVM_LIBC_FUNCTION(double, atan, (double x)) { return __nv_atan(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/atan2.cpp b/libc/src/math/nvptx/atan2.cpp
similarity index 91%
rename from libc/src/math/gpu/vendor/atan2.cpp
rename to libc/src/math/nvptx/atan2.cpp
index 94e215e52ca602..0c54e0e0489964 100644
--- a/libc/src/math/gpu/vendor/atan2.cpp
+++ b/libc/src/math/nvptx/atan2.cpp
@@ -9,12 +9,12 @@
#include "src/math/atan2.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(double, atan2, (double x, double y)) {
- return internal::atan2(x, y);
+ return __nv_atan2(x, y);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/atan2f.cpp b/libc/src/math/nvptx/atan2f.cpp
similarity index 91%
rename from libc/src/math/gpu/vendor/atan2f.cpp
rename to libc/src/math/nvptx/atan2f.cpp
index 70caa568e32d93..c3327d92c97e9a 100644
--- a/libc/src/math/gpu/vendor/atan2f.cpp
+++ b/libc/src/math/nvptx/atan2f.cpp
@@ -9,12 +9,12 @@
#include "src/math/atan2f.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(float, atan2f, (float x, float y)) {
- return internal::atan2f(x, y);
+ return __nv_atan2f(x, y);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/atanf.cpp b/libc/src/math/nvptx/atanf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/atanf.cpp
rename to libc/src/math/nvptx/atanf.cpp
index 132c43d9e3af7c..5595262977323a 100644
--- a/libc/src/math/gpu/vendor/atanf.cpp
+++ b/libc/src/math/nvptx/atanf.cpp
@@ -9,10 +9,10 @@
#include "src/math/atanf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, atanf, (float x)) { return internal::atanf(x); }
+LLVM_LIBC_FUNCTION(float, atanf, (float x)) { return __nv_atanf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/atanh.cpp b/libc/src/math/nvptx/atanh.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/atanh.cpp
rename to libc/src/math/nvptx/atanh.cpp
index 07a75fcbbfc7c6..6699d959df187a 100644
--- a/libc/src/math/gpu/vendor/atanh.cpp
+++ b/libc/src/math/nvptx/atanh.cpp
@@ -9,10 +9,10 @@
#include "src/math/atanh.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, atanh, (double x)) { return internal::atanh(x); }
+LLVM_LIBC_FUNCTION(double, atanh, (double x)) { return __nv_atanh(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/atanhf.cpp b/libc/src/math/nvptx/atanhf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/atanhf.cpp
rename to libc/src/math/nvptx/atanhf.cpp
index 521c4133243d98..526b7b3e371225 100644
--- a/libc/src/math/gpu/vendor/atanhf.cpp
+++ b/libc/src/math/nvptx/atanhf.cpp
@@ -9,10 +9,10 @@
#include "src/math/atanhf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, atanhf, (float x)) { return internal::atanhf(x); }
+LLVM_LIBC_FUNCTION(float, atanhf, (float x)) { return __nv_atanhf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/ceil.cpp b/libc/src/math/nvptx/ceil.cpp
new file mode 100644
index 00000000000000..ad1407d61f620a
--- /dev/null
+++ b/libc/src/math/nvptx/ceil.cpp
@@ -0,0 +1,16 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/ceil.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, ceil, (double x)) { return __builtin_ceil(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/ceilf.cpp b/libc/src/math/nvptx/ceilf.cpp
new file mode 100644
index 00000000000000..c4fc58d9360385
--- /dev/null
+++ b/libc/src/math/nvptx/ceilf.cpp
@@ -0,0 +1,16 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/ceilf.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, ceilf, (float x)) { return __builtin_ceilf(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/copysign.cpp b/libc/src/math/nvptx/copysign.cpp
new file mode 100644
index 00000000000000..6f804bdb90a1f0
--- /dev/null
+++ b/libc/src/math/nvptx/copysign.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/copysign.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, copysign, (double x, double y)) {
+ return __builtin_copysign(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/copysignf.cpp b/libc/src/math/nvptx/copysignf.cpp
new file mode 100644
index 00000000000000..4d7e132462ac90
--- /dev/null
+++ b/libc/src/math/nvptx/copysignf.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/copysignf.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, copysignf, (float x, float y)) {
+ return __builtin_copysignf(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/cos.cpp b/libc/src/math/nvptx/cos.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/cos.cpp
rename to libc/src/math/nvptx/cos.cpp
index 37c7507911ec2e..185ad3cf921594 100644
--- a/libc/src/math/gpu/vendor/cos.cpp
+++ b/libc/src/math/nvptx/cos.cpp
@@ -9,10 +9,10 @@
#include "src/math/cos.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, cos, (double x)) { return internal::cos(x); }
+LLVM_LIBC_FUNCTION(double, cos, (double x)) { return __nv_cos(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/cosf.cpp b/libc/src/math/nvptx/cosf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/cosf.cpp
rename to libc/src/math/nvptx/cosf.cpp
index 1bd42ba37e92d9..3d34de4be51bc4 100644
--- a/libc/src/math/gpu/vendor/cosf.cpp
+++ b/libc/src/math/nvptx/cosf.cpp
@@ -9,10 +9,10 @@
#include "src/math/cosf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, cosf, (float x)) { return internal::cosf(x); }
+LLVM_LIBC_FUNCTION(float, cosf, (float x)) { return __nv_cosf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/cosh.cpp b/libc/src/math/nvptx/cosh.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/cosh.cpp
rename to libc/src/math/nvptx/cosh.cpp
index 3be05e58b51d96..179864c5f910d7 100644
--- a/libc/src/math/gpu/vendor/cosh.cpp
+++ b/libc/src/math/nvptx/cosh.cpp
@@ -9,10 +9,10 @@
#include "src/math/cosh.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, cosh, (double x)) { return internal::cosh(x); }
+LLVM_LIBC_FUNCTION(double, cosh, (double x)) { return __nv_cosh(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/coshf.cpp b/libc/src/math/nvptx/coshf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/coshf.cpp
rename to libc/src/math/nvptx/coshf.cpp
index 1b945bbd7c47e4..9147499db97cff 100644
--- a/libc/src/math/gpu/vendor/coshf.cpp
+++ b/libc/src/math/nvptx/coshf.cpp
@@ -9,10 +9,10 @@
#include "src/math/coshf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, coshf, (float x)) { return internal::coshf(x); }
+LLVM_LIBC_FUNCTION(float, coshf, (float x)) { return __nv_coshf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/nvptx/declarations.h b/libc/src/math/nvptx/declarations.h
similarity index 100%
rename from libc/src/math/gpu/vendor/nvptx/declarations.h
rename to libc/src/math/nvptx/declarations.h
diff --git a/libc/src/math/gpu/vendor/erf.cpp b/libc/src/math/nvptx/erf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/erf.cpp
rename to libc/src/math/nvptx/erf.cpp
index 190321ca25992e..5ea0177d5cd346 100644
--- a/libc/src/math/gpu/vendor/erf.cpp
+++ b/libc/src/math/nvptx/erf.cpp
@@ -9,10 +9,10 @@
#include "src/math/erf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, erf, (double x)) { return internal::erf(x); }
+LLVM_LIBC_FUNCTION(double, erf, (double x)) { return __nv_erf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/erff.cpp b/libc/src/math/nvptx/erff.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/erff.cpp
rename to libc/src/math/nvptx/erff.cpp
index a5a08be54b0756..03fdceace8e9e0 100644
--- a/libc/src/math/gpu/vendor/erff.cpp
+++ b/libc/src/math/nvptx/erff.cpp
@@ -9,10 +9,10 @@
#include "src/math/erff.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, erff, (float x)) { return internal::erff(x); }
+LLVM_LIBC_FUNCTION(float, erff, (float x)) { return __nv_erff(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/exp.cpp b/libc/src/math/nvptx/exp.cpp
new file mode 100644
index 00000000000000..6bbe87ba2e7836
--- /dev/null
+++ b/libc/src/math/nvptx/exp.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU exp function ----------------------------===//
+//
+// 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/exp.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, exp, (double x)) { return __nv_exp(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/exp10.cpp b/libc/src/math/nvptx/exp10.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/exp10.cpp
rename to libc/src/math/nvptx/exp10.cpp
index 8557a33f018844..11bb734fd11319 100644
--- a/libc/src/math/gpu/vendor/exp10.cpp
+++ b/libc/src/math/nvptx/exp10.cpp
@@ -9,10 +9,10 @@
#include "src/math/exp10.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, exp10, (double x)) { return internal::exp10(x); }
+LLVM_LIBC_FUNCTION(double, exp10, (double x)) { return __nv_exp10(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/exp10f.cpp b/libc/src/math/nvptx/exp10f.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/exp10f.cpp
rename to libc/src/math/nvptx/exp10f.cpp
index 844809355087b4..4e3121a0b46e23 100644
--- a/libc/src/math/gpu/vendor/exp10f.cpp
+++ b/libc/src/math/nvptx/exp10f.cpp
@@ -9,10 +9,10 @@
#include "src/math/exp10f.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, exp10f, (float x)) { return internal::exp10f(x); }
+LLVM_LIBC_FUNCTION(float, exp10f, (float x)) { return __nv_exp10f(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/exp2.cpp b/libc/src/math/nvptx/exp2.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/exp2.cpp
rename to libc/src/math/nvptx/exp2.cpp
index ffa23d810a9b1e..35fc27b3a26a66 100644
--- a/libc/src/math/gpu/vendor/exp2.cpp
+++ b/libc/src/math/nvptx/exp2.cpp
@@ -9,10 +9,10 @@
#include "src/math/exp2.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, exp2, (double x)) { return internal::exp2(x); }
+LLVM_LIBC_FUNCTION(double, exp2, (double x)) { return __nv_exp2(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/exp2f.cpp b/libc/src/math/nvptx/exp2f.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/exp2f.cpp
rename to libc/src/math/nvptx/exp2f.cpp
index cb61557383dfa7..8d137346fe00e9 100644
--- a/libc/src/math/gpu/vendor/exp2f.cpp
+++ b/libc/src/math/nvptx/exp2f.cpp
@@ -9,10 +9,10 @@
#include "src/math/exp2f.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, exp2f, (float x)) { return internal::exp2f(x); }
+LLVM_LIBC_FUNCTION(float, exp2f, (float x)) { return __nv_exp2f(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/expf.cpp b/libc/src/math/nvptx/expf.cpp
new file mode 100644
index 00000000000000..a6362bd7346113
--- /dev/null
+++ b/libc/src/math/nvptx/expf.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/expf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, expf, (float x)) { return __nv_expf(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/expm1.cpp b/libc/src/math/nvptx/expm1.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/expm1.cpp
rename to libc/src/math/nvptx/expm1.cpp
index 6ac5f753b9e4df..0331903b8fd869 100644
--- a/libc/src/math/gpu/vendor/expm1.cpp
+++ b/libc/src/math/nvptx/expm1.cpp
@@ -9,10 +9,10 @@
#include "src/math/expm1.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, expm1, (double x)) { return internal::expm1(x); }
+LLVM_LIBC_FUNCTION(double, expm1, (double x)) { return __nv_expm1(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/expm1f.cpp b/libc/src/math/nvptx/expm1f.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/expm1f.cpp
rename to libc/src/math/nvptx/expm1f.cpp
index c5497797dbe1b3..7b74c548f3df61 100644
--- a/libc/src/math/gpu/vendor/expm1f.cpp
+++ b/libc/src/math/nvptx/expm1f.cpp
@@ -9,10 +9,10 @@
#include "src/math/expm1f.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, expm1f, (float x)) { return internal::expm1f(x); }
+LLVM_LIBC_FUNCTION(float, expm1f, (float x)) { return __nv_expm1f(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/fabs.cpp b/libc/src/math/nvptx/fabs.cpp
new file mode 100644
index 00000000000000..c0d063d50ae53d
--- /dev/null
+++ b/libc/src/math/nvptx/fabs.cpp
@@ -0,0 +1,16 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/fabs.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, fabs, (double x)) { return __builtin_fabs(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/fabsf.cpp b/libc/src/math/nvptx/fabsf.cpp
new file mode 100644
index 00000000000000..398ffd0c74c074
--- /dev/null
+++ b/libc/src/math/nvptx/fabsf.cpp
@@ -0,0 +1,16 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/fabsf.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, fabsf, (float x)) { return __builtin_fabsf(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/fdim.cpp b/libc/src/math/nvptx/fdim.cpp
similarity index 91%
rename from libc/src/math/gpu/vendor/fdim.cpp
rename to libc/src/math/nvptx/fdim.cpp
index f30dafb46e5488..2f1ff518002662 100644
--- a/libc/src/math/gpu/vendor/fdim.cpp
+++ b/libc/src/math/nvptx/fdim.cpp
@@ -9,12 +9,12 @@
#include "src/math/fdim.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(double, fdim, (double x, double y)) {
- return internal::fdim(x, y);
+ return __nv_fdim(x, y);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/fdimf.cpp b/libc/src/math/nvptx/fdimf.cpp
similarity index 91%
rename from libc/src/math/gpu/vendor/fdimf.cpp
rename to libc/src/math/nvptx/fdimf.cpp
index e30736206a9a92..c24e6be72482aa 100644
--- a/libc/src/math/gpu/vendor/fdimf.cpp
+++ b/libc/src/math/nvptx/fdimf.cpp
@@ -9,12 +9,12 @@
#include "src/math/fdimf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(float, fdimf, (float x, float y)) {
- return internal::fdimf(x, y);
+ return __nv_fdimf(x, y);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/floor.cpp b/libc/src/math/nvptx/floor.cpp
new file mode 100644
index 00000000000000..eada89c178d75c
--- /dev/null
+++ b/libc/src/math/nvptx/floor.cpp
@@ -0,0 +1,16 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/floor.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, floor, (double x)) { return __builtin_floor(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/floorf.cpp b/libc/src/math/nvptx/floorf.cpp
new file mode 100644
index 00000000000000..a5611c515a88d4
--- /dev/null
+++ b/libc/src/math/nvptx/floorf.cpp
@@ -0,0 +1,16 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/floorf.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, floorf, (float x)) { return __builtin_floorf(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/fma.cpp b/libc/src/math/nvptx/fma.cpp
new file mode 100644
index 00000000000000..41a6ddf60dbc35
--- /dev/null
+++ b/libc/src/math/nvptx/fma.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/fma.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, fma, (double x, double y, double z)) {
+ return __builtin_fma(x, y, z);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/fmaf.cpp b/libc/src/math/nvptx/fmaf.cpp
new file mode 100644
index 00000000000000..c948e32f77eb92
--- /dev/null
+++ b/libc/src/math/nvptx/fmaf.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/fmaf.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, fmaf, (float x, float y, float z)) {
+ return __builtin_fmaf(x, y, z);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/fmax.cpp b/libc/src/math/nvptx/fmax.cpp
new file mode 100644
index 00000000000000..a2c35371d12b6a
--- /dev/null
+++ b/libc/src/math/nvptx/fmax.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/fmax.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, fmax, (double x, double y)) {
+ return __builtin_fmax(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/fmaxf.cpp b/libc/src/math/nvptx/fmaxf.cpp
new file mode 100644
index 00000000000000..67178b3e273575
--- /dev/null
+++ b/libc/src/math/nvptx/fmaxf.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/fmaxf.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, fmaxf, (float x, float y)) {
+ return __builtin_fmaxf(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/fmin.cpp b/libc/src/math/nvptx/fmin.cpp
new file mode 100644
index 00000000000000..7303adcd347ee9
--- /dev/null
+++ b/libc/src/math/nvptx/fmin.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/fmin.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, fmin, (double x, double y)) {
+ return __builtin_fmin(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/fminf.cpp b/libc/src/math/nvptx/fminf.cpp
new file mode 100644
index 00000000000000..bbf0c677b5e3ae
--- /dev/null
+++ b/libc/src/math/nvptx/fminf.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/fminf.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, fminf, (float x, float y)) {
+ return __builtin_fminf(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/fmod.cpp b/libc/src/math/nvptx/fmod.cpp
new file mode 100644
index 00000000000000..0654cdd2abe081
--- /dev/null
+++ b/libc/src/math/nvptx/fmod.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/fmod.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, fmod, (double x, double y)) {
+ return __builtin_fmod(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/fmodf.cpp b/libc/src/math/nvptx/fmodf.cpp
new file mode 100644
index 00000000000000..b689046468fbe5
--- /dev/null
+++ b/libc/src/math/nvptx/fmodf.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/fmodf.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, fmodf, (float x, float y)) {
+ return __builtin_fmodf(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/frexp.cpp b/libc/src/math/nvptx/frexp.cpp
similarity index 91%
rename from libc/src/math/gpu/vendor/frexp.cpp
rename to libc/src/math/nvptx/frexp.cpp
index 5fc2c1409c6e93..2423961f7c61ea 100644
--- a/libc/src/math/gpu/vendor/frexp.cpp
+++ b/libc/src/math/nvptx/frexp.cpp
@@ -9,12 +9,12 @@
#include "src/math/frexp.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(double, frexp, (double x, int *p)) {
- return internal::frexp(x, p);
+ return __nv_frexp(x, p);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/frexpf.cpp b/libc/src/math/nvptx/frexpf.cpp
similarity index 91%
rename from libc/src/math/gpu/vendor/frexpf.cpp
rename to libc/src/math/nvptx/frexpf.cpp
index e928d375e03db3..f1ea29068777bc 100644
--- a/libc/src/math/gpu/vendor/frexpf.cpp
+++ b/libc/src/math/nvptx/frexpf.cpp
@@ -9,12 +9,12 @@
#include "src/math/frexpf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(float, frexpf, (float x, int *p)) {
- return internal::frexpf(x, p);
+ return __nv_frexpf(x, p);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/hypot.cpp b/libc/src/math/nvptx/hypot.cpp
similarity index 91%
rename from libc/src/math/gpu/vendor/hypot.cpp
rename to libc/src/math/nvptx/hypot.cpp
index 45b629e3e28585..28bf04aa4a2bda 100644
--- a/libc/src/math/gpu/vendor/hypot.cpp
+++ b/libc/src/math/nvptx/hypot.cpp
@@ -9,12 +9,12 @@
#include "src/math/hypot.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(double, hypot, (double x, double y)) {
- return internal::hypot(x, y);
+ return __nv_hypot(x, y);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/hypotf.cpp b/libc/src/math/nvptx/hypotf.cpp
similarity index 91%
rename from libc/src/math/gpu/vendor/hypotf.cpp
rename to libc/src/math/nvptx/hypotf.cpp
index 533e9dcb8dbfb0..c506aab1acc0b4 100644
--- a/libc/src/math/gpu/vendor/hypotf.cpp
+++ b/libc/src/math/nvptx/hypotf.cpp
@@ -9,12 +9,12 @@
#include "src/math/hypotf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(float, hypotf, (float x, float y)) {
- return internal::hypotf(x, y);
+ return __nv_hypotf(x, y);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/ilogb.cpp b/libc/src/math/nvptx/ilogb.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/ilogb.cpp
rename to libc/src/math/nvptx/ilogb.cpp
index 1d075027b41c05..fc75e2fd847a86 100644
--- a/libc/src/math/gpu/vendor/ilogb.cpp
+++ b/libc/src/math/nvptx/ilogb.cpp
@@ -9,10 +9,10 @@
#include "src/math/ilogb.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(int, ilogb, (double x)) { return internal::ilogb(x); }
+LLVM_LIBC_FUNCTION(int, ilogb, (double x)) { return __nv_ilogb(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/ilogbf.cpp b/libc/src/math/nvptx/ilogbf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/ilogbf.cpp
rename to libc/src/math/nvptx/ilogbf.cpp
index 8dc2ff0a374af2..3d14fcfa878fbf 100644
--- a/libc/src/math/gpu/vendor/ilogbf.cpp
+++ b/libc/src/math/nvptx/ilogbf.cpp
@@ -9,10 +9,10 @@
#include "src/math/ilogbf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(int, ilogbf, (float x)) { return internal::ilogbf(x); }
+LLVM_LIBC_FUNCTION(int, ilogbf, (float x)) { return __nv_ilogbf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/ldexp.cpp b/libc/src/math/nvptx/ldexp.cpp
similarity index 91%
rename from libc/src/math/gpu/vendor/ldexp.cpp
rename to libc/src/math/nvptx/ldexp.cpp
index f760a42563305b..761dc4816b7a4f 100644
--- a/libc/src/math/gpu/vendor/ldexp.cpp
+++ b/libc/src/math/nvptx/ldexp.cpp
@@ -9,12 +9,12 @@
#include "src/math/ldexp.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(double, ldexp, (double x, int y)) {
- return internal::ldexp(x, y);
+ return __nv_ldexp(x, y);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/ldexpf.cpp b/libc/src/math/nvptx/ldexpf.cpp
similarity index 91%
rename from libc/src/math/gpu/vendor/ldexpf.cpp
rename to libc/src/math/nvptx/ldexpf.cpp
index d00d39115ffc33..2d4c556a271406 100644
--- a/libc/src/math/gpu/vendor/ldexpf.cpp
+++ b/libc/src/math/nvptx/ldexpf.cpp
@@ -9,12 +9,12 @@
#include "src/math/ldexpf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(float, ldexpf, (float x, int y)) {
- return internal::ldexpf(x, y);
+ return __nv_ldexpf(x, y);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/llrint.cpp b/libc/src/math/nvptx/llrint.cpp
new file mode 100644
index 00000000000000..8f95e75e779b51
--- /dev/null
+++ b/libc/src/math/nvptx/llrint.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/llrint.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(long long, llrint, (double x)) { return __nv_llrint(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/llrintf.cpp b/libc/src/math/nvptx/llrintf.cpp
new file mode 100644
index 00000000000000..1432ffbd1bdad3
--- /dev/null
+++ b/libc/src/math/nvptx/llrintf.cpp
@@ -0,0 +1,18 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/llrintf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(long long, llrintf, (float x)) { return __nv_llrintf(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/llround.cpp b/libc/src/math/nvptx/llround.cpp
new file mode 100644
index 00000000000000..afd98308730a62
--- /dev/null
+++ b/libc/src/math/nvptx/llround.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU llround function ------------------------===//
+//
+// 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/llround.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(long long, llround, (double x)) {
+ return __builtin_llround(x);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/llroundf.cpp b/libc/src/math/nvptx/llroundf.cpp
new file mode 100644
index 00000000000000..897ed15b692848
--- /dev/null
+++ b/libc/src/math/nvptx/llroundf.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU llroundf function -----------------------===//
+//
+// 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/llroundf.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(long long, llroundf, (float x)) {
+ return __builtin_lroundf(x);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/log.cpp b/libc/src/math/nvptx/log.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/log.cpp
rename to libc/src/math/nvptx/log.cpp
index a97689abcc2171..26b6dfa607b94d 100644
--- a/libc/src/math/gpu/vendor/log.cpp
+++ b/libc/src/math/nvptx/log.cpp
@@ -9,10 +9,10 @@
#include "src/math/log.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, log, (double x)) { return internal::log(x); }
+LLVM_LIBC_FUNCTION(double, log, (double x)) { return __nv_log(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/log10.cpp b/libc/src/math/nvptx/log10.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/log10.cpp
rename to libc/src/math/nvptx/log10.cpp
index c7a917a75cc969..ff27025395672b 100644
--- a/libc/src/math/gpu/vendor/log10.cpp
+++ b/libc/src/math/nvptx/log10.cpp
@@ -9,10 +9,10 @@
#include "src/math/log10.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, log10, (double x)) { return internal::log10(x); }
+LLVM_LIBC_FUNCTION(double, log10, (double x)) { return __nv_log10(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/log10f.cpp b/libc/src/math/nvptx/log10f.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/log10f.cpp
rename to libc/src/math/nvptx/log10f.cpp
index 489f5f558be1f5..af903b60a7ce68 100644
--- a/libc/src/math/gpu/vendor/log10f.cpp
+++ b/libc/src/math/nvptx/log10f.cpp
@@ -9,10 +9,10 @@
#include "src/math/log10f.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, log10f, (float x)) { return internal::log10f(x); }
+LLVM_LIBC_FUNCTION(float, log10f, (float x)) { return __nv_log10f(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/log1p.cpp b/libc/src/math/nvptx/log1p.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/log1p.cpp
rename to libc/src/math/nvptx/log1p.cpp
index 720d23e2f952b8..47bc96b0d8812c 100644
--- a/libc/src/math/gpu/vendor/log1p.cpp
+++ b/libc/src/math/nvptx/log1p.cpp
@@ -9,10 +9,10 @@
#include "src/math/log1p.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, log1p, (double x)) { return internal::log1p(x); }
+LLVM_LIBC_FUNCTION(double, log1p, (double x)) { return __nv_log1p(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/log1pf.cpp b/libc/src/math/nvptx/log1pf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/log1pf.cpp
rename to libc/src/math/nvptx/log1pf.cpp
index 96ad48b529cf6f..bfa4f7f22d54cd 100644
--- a/libc/src/math/gpu/vendor/log1pf.cpp
+++ b/libc/src/math/nvptx/log1pf.cpp
@@ -9,10 +9,10 @@
#include "src/math/log1pf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, log1pf, (float x)) { return internal::log1pf(x); }
+LLVM_LIBC_FUNCTION(float, log1pf, (float x)) { return __nv_log1pf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/log2.cpp b/libc/src/math/nvptx/log2.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/log2.cpp
rename to libc/src/math/nvptx/log2.cpp
index 9fc8a81e7e7570..86a980de65d4d7 100644
--- a/libc/src/math/gpu/vendor/log2.cpp
+++ b/libc/src/math/nvptx/log2.cpp
@@ -9,10 +9,10 @@
#include "src/math/log2.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, log2, (double x)) { return internal::log2(x); }
+LLVM_LIBC_FUNCTION(double, log2, (double x)) { return __nv_log2(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/log2f.cpp b/libc/src/math/nvptx/log2f.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/log2f.cpp
rename to libc/src/math/nvptx/log2f.cpp
index 62df41b69b0bd8..5ce46291610dd3 100644
--- a/libc/src/math/gpu/vendor/log2f.cpp
+++ b/libc/src/math/nvptx/log2f.cpp
@@ -9,10 +9,10 @@
#include "src/math/log2f.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, log2f, (float x)) { return internal::log2f(x); }
+LLVM_LIBC_FUNCTION(float, log2f, (float x)) { return __nv_log2f(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/logb.cpp b/libc/src/math/nvptx/logb.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/logb.cpp
rename to libc/src/math/nvptx/logb.cpp
index 5dea57d41b08b1..b620b16184fc90 100644
--- a/libc/src/math/gpu/vendor/logb.cpp
+++ b/libc/src/math/nvptx/logb.cpp
@@ -9,10 +9,10 @@
#include "src/math/logb.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, logb, (double x)) { return internal::logb(x); }
+LLVM_LIBC_FUNCTION(double, logb, (double x)) { return __nv_logb(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/logbf.cpp b/libc/src/math/nvptx/logbf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/logbf.cpp
rename to libc/src/math/nvptx/logbf.cpp
index 1a59df3e09a89b..f19f0320db9d48 100644
--- a/libc/src/math/gpu/vendor/logbf.cpp
+++ b/libc/src/math/nvptx/logbf.cpp
@@ -9,10 +9,10 @@
#include "src/math/logbf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, logbf, (float x)) { return internal::logbf(x); }
+LLVM_LIBC_FUNCTION(float, logbf, (float x)) { return __nv_logbf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/logf.cpp b/libc/src/math/nvptx/logf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/logf.cpp
rename to libc/src/math/nvptx/logf.cpp
index 527b028e100d5e..6deb482c0ace39 100644
--- a/libc/src/math/gpu/vendor/logf.cpp
+++ b/libc/src/math/nvptx/logf.cpp
@@ -9,10 +9,10 @@
#include "src/math/logf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, logf, (float x)) { return internal::logf(x); }
+LLVM_LIBC_FUNCTION(float, logf, (float x)) { return __nv_logf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/lrint.cpp b/libc/src/math/nvptx/lrint.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/lrint.cpp
rename to libc/src/math/nvptx/lrint.cpp
index a08996b755b528..8585f4ce53a4d1 100644
--- a/libc/src/math/gpu/vendor/lrint.cpp
+++ b/libc/src/math/nvptx/lrint.cpp
@@ -9,10 +9,10 @@
#include "src/math/lrint.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(long, lrint, (double x)) { return internal::lrint(x); }
+LLVM_LIBC_FUNCTION(long, lrint, (double x)) { return __nv_lrint(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/lrintf.cpp b/libc/src/math/nvptx/lrintf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/lrintf.cpp
rename to libc/src/math/nvptx/lrintf.cpp
index 695a9b8202cf4a..312a9469fade01 100644
--- a/libc/src/math/gpu/vendor/lrintf.cpp
+++ b/libc/src/math/nvptx/lrintf.cpp
@@ -9,10 +9,10 @@
#include "src/math/lrintf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(long, lrintf, (float x)) { return internal::lrintf(x); }
+LLVM_LIBC_FUNCTION(long, lrintf, (float x)) { return __nv_lrintf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/lround.cpp b/libc/src/math/nvptx/lround.cpp
new file mode 100644
index 00000000000000..51e8f2245af8ea
--- /dev/null
+++ b/libc/src/math/nvptx/lround.cpp
@@ -0,0 +1,16 @@
+//===-- Implementation of the GPU lround function -------------------------===//
+//
+// 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/lround.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(long, lround, (double x)) { return __builtin_lround(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/lroundf.cpp b/libc/src/math/nvptx/lroundf.cpp
new file mode 100644
index 00000000000000..2a6fe7200d8cba
--- /dev/null
+++ b/libc/src/math/nvptx/lroundf.cpp
@@ -0,0 +1,16 @@
+//===-- Implementation of the GPU lroundf function ------------------------===//
+//
+// 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/lroundf.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(long, lroundf, (float x)) { return __builtin_lroundf(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/modf.cpp b/libc/src/math/nvptx/modf.cpp
new file mode 100644
index 00000000000000..07dbbd6059c35f
--- /dev/null
+++ b/libc/src/math/nvptx/modf.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU modf function ---------------------------===//
+//
+// 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/modf.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, modf, (double x, double *iptr)) {
+ return __builtin_modf(x, iptr);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/modff.cpp b/libc/src/math/nvptx/modff.cpp
new file mode 100644
index 00000000000000..ad35f9006b5122
--- /dev/null
+++ b/libc/src/math/nvptx/modff.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU modff function --------------------------===//
+//
+// 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/modff.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, modff, (float x, float *iptr)) {
+ return __builtin_modff(x, iptr);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/nearbyint.cpp b/libc/src/math/nvptx/nearbyint.cpp
new file mode 100644
index 00000000000000..9c7b600df7082e
--- /dev/null
+++ b/libc/src/math/nvptx/nearbyint.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU nearbyint function ----------------------===//
+//
+// 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/nearbyint.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, nearbyint, (double x)) {
+ return __builtin_nearbyint(x);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/nearbyintf.cpp b/libc/src/math/nvptx/nearbyintf.cpp
new file mode 100644
index 00000000000000..7fbe9f4f0e0bee
--- /dev/null
+++ b/libc/src/math/nvptx/nearbyintf.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU nearbyintf function ---------------------===//
+//
+// 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/nearbyintf.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, nearbyintf, (float x)) {
+ return __builtin_nearbyintf(x);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/nextafter.cpp b/libc/src/math/nvptx/nextafter.cpp
similarity index 90%
rename from libc/src/math/gpu/vendor/nextafter.cpp
rename to libc/src/math/nvptx/nextafter.cpp
index f88e17f1090842..171aaad6f7cc17 100644
--- a/libc/src/math/gpu/vendor/nextafter.cpp
+++ b/libc/src/math/nvptx/nextafter.cpp
@@ -9,12 +9,12 @@
#include "src/math/nextafter.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(double, nextafter, (double x, double y)) {
- return internal::nextafter(x, y);
+ return __nv_nextafter(x, y);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/nextafterf.cpp b/libc/src/math/nvptx/nextafterf.cpp
similarity index 90%
rename from libc/src/math/gpu/vendor/nextafterf.cpp
rename to libc/src/math/nvptx/nextafterf.cpp
index 7a39dc8fc814ea..a45937c0dc8ba5 100644
--- a/libc/src/math/gpu/vendor/nextafterf.cpp
+++ b/libc/src/math/nvptx/nextafterf.cpp
@@ -9,12 +9,12 @@
#include "src/math/nextafterf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(float, nextafterf, (float x, float y)) {
- return internal::nextafterf(x, y);
+ return __nv_nextafterf(x, y);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/nvptx/nvptx.h b/libc/src/math/nvptx/nvptx.h
similarity index 100%
rename from libc/src/math/gpu/vendor/nvptx/nvptx.h
rename to libc/src/math/nvptx/nvptx.h
diff --git a/libc/src/math/nvptx/pow.cpp b/libc/src/math/nvptx/pow.cpp
new file mode 100644
index 00000000000000..7de3c9e7e54485
--- /dev/null
+++ b/libc/src/math/nvptx/pow.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the pow 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/pow.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, pow, (double x, double y)) { return __nv_pow(x, y); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/powf.cpp b/libc/src/math/nvptx/powf.cpp
new file mode 100644
index 00000000000000..f9f7dbae63ac54
--- /dev/null
+++ b/libc/src/math/nvptx/powf.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the powf 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/powf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, powf, (float x, float y)) { return __nv_powf(x, y); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/remainder.cpp b/libc/src/math/nvptx/remainder.cpp
new file mode 100644
index 00000000000000..89b235f9c22af3
--- /dev/null
+++ b/libc/src/math/nvptx/remainder.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU remainder function ----------------------===//
+//
+// 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/remainder.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, remainder, (double x, double y)) {
+ return __builtin_remainder(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/remainderf.cpp b/libc/src/math/nvptx/remainderf.cpp
new file mode 100644
index 00000000000000..9fee6f856dc8b7
--- /dev/null
+++ b/libc/src/math/nvptx/remainderf.cpp
@@ -0,0 +1,18 @@
+//===-- Implementation of the GPU remainderf function ---------------------===//
+//
+// 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/remainderf.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, remainderf, (float x, float y)) {
+ return __builtin_remainderf(x, y);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/remquo.cpp b/libc/src/math/nvptx/remquo.cpp
similarity index 90%
rename from libc/src/math/gpu/vendor/remquo.cpp
rename to libc/src/math/nvptx/remquo.cpp
index e92c9b3c2a6e3b..da69a20f8f4f84 100644
--- a/libc/src/math/gpu/vendor/remquo.cpp
+++ b/libc/src/math/nvptx/remquo.cpp
@@ -9,12 +9,12 @@
#include "src/math/remquo.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(double, remquo, (double x, double y, int *quo)) {
- return internal::remquo(x, y, quo);
+ return __nv_remquo(x, y, quo);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/remquof.cpp b/libc/src/math/nvptx/remquof.cpp
similarity index 90%
rename from libc/src/math/gpu/vendor/remquof.cpp
rename to libc/src/math/nvptx/remquof.cpp
index b234885aa88c03..dcfba5d7b5fa94 100644
--- a/libc/src/math/gpu/vendor/remquof.cpp
+++ b/libc/src/math/nvptx/remquof.cpp
@@ -9,12 +9,12 @@
#include "src/math/remquof.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(float, remquof, (float x, float y, int *quo)) {
- return internal::remquof(x, y, quo);
+ return __nv_remquof(x, y, quo);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/rint.cpp b/libc/src/math/nvptx/rint.cpp
new file mode 100644
index 00000000000000..44d494a8ed57af
--- /dev/null
+++ b/libc/src/math/nvptx/rint.cpp
@@ -0,0 +1,16 @@
+//===-- Implementation of the GPU rint function ---------------------------===//
+//
+// 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/rint.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, rint, (double x)) { return __builtin_rint(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/rintf.cpp b/libc/src/math/nvptx/rintf.cpp
new file mode 100644
index 00000000000000..daf98d94360511
--- /dev/null
+++ b/libc/src/math/nvptx/rintf.cpp
@@ -0,0 +1,16 @@
+//===-- Implementation of the GPU rintf function --------------------------===//
+//
+// 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/rintf.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, rintf, (float x)) { return __builtin_rintf(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/round.cpp b/libc/src/math/nvptx/round.cpp
new file mode 100644
index 00000000000000..9d8b5582f0407a
--- /dev/null
+++ b/libc/src/math/nvptx/round.cpp
@@ -0,0 +1,16 @@
+//===-- Implementation of the GPU round function --------------------------===//
+//
+// 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/round.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, round, (double x)) { return __builtin_round(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/roundf.cpp b/libc/src/math/nvptx/roundf.cpp
new file mode 100644
index 00000000000000..8743e4eb7fb8d7
--- /dev/null
+++ b/libc/src/math/nvptx/roundf.cpp
@@ -0,0 +1,16 @@
+//===-- Implementation of the GPU roundf function -------------------------===//
+//
+// 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/roundf.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, roundf, (float x)) { return __builtin_roundf(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/scalbn.cpp b/libc/src/math/nvptx/scalbn.cpp
similarity index 91%
rename from libc/src/math/gpu/vendor/scalbn.cpp
rename to libc/src/math/nvptx/scalbn.cpp
index 435533a90501b7..80374db4c1c93d 100644
--- a/libc/src/math/gpu/vendor/scalbn.cpp
+++ b/libc/src/math/nvptx/scalbn.cpp
@@ -9,12 +9,12 @@
#include "src/math/scalbn.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(double, scalbn, (double x, int y)) {
- return internal::scalbn(x, y);
+ return __nv_scalbn(x, y);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/scalbnf.cpp b/libc/src/math/nvptx/scalbnf.cpp
similarity index 90%
rename from libc/src/math/gpu/vendor/scalbnf.cpp
rename to libc/src/math/nvptx/scalbnf.cpp
index 0a4844c8a433ef..24fa3a5ed698ae 100644
--- a/libc/src/math/gpu/vendor/scalbnf.cpp
+++ b/libc/src/math/nvptx/scalbnf.cpp
@@ -9,12 +9,12 @@
#include "src/math/scalbnf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(float, scalbnf, (float x, int y)) {
- return internal::scalbnf(x, y);
+ return __nv_scalbnf(x, y);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/sin.cpp b/libc/src/math/nvptx/sin.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/sin.cpp
rename to libc/src/math/nvptx/sin.cpp
index 96e07c99b71a33..1bff129a0151c0 100644
--- a/libc/src/math/gpu/vendor/sin.cpp
+++ b/libc/src/math/nvptx/sin.cpp
@@ -9,10 +9,10 @@
#include "src/math/sin.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, sin, (double x)) { return internal::sin(x); }
+LLVM_LIBC_FUNCTION(double, sin, (double x)) { return __nv_sin(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/sincos.cpp b/libc/src/math/nvptx/sincos.cpp
similarity index 89%
rename from libc/src/math/gpu/vendor/sincos.cpp
rename to libc/src/math/nvptx/sincos.cpp
index d882157b0bc8d7..73f92cfb7c3488 100644
--- a/libc/src/math/gpu/vendor/sincos.cpp
+++ b/libc/src/math/nvptx/sincos.cpp
@@ -9,12 +9,12 @@
#include "src/math/sincos.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
LLVM_LIBC_FUNCTION(void, sincos, (double x, double *sinptr, double *cosptr)) {
- return internal::sincos(x, sinptr, cosptr);
+ return __nv_sincos(x, sinptr, cosptr);
}
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/sincosf.cpp b/libc/src/math/nvptx/sincosf.cpp
new file mode 100644
index 00000000000000..d053aa38151b63
--- /dev/null
+++ b/libc/src/math/nvptx/sincosf.cpp
@@ -0,0 +1,20 @@
+//===-- Implementation of the sincosf 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/sincosf.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(void, sincosf, (float x, float *sinptr, float *cosptr)) {
+ return __nv_sincosf(x, sinptr, cosptr);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/sinf.cpp b/libc/src/math/nvptx/sinf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/sinf.cpp
rename to libc/src/math/nvptx/sinf.cpp
index af93227ab63bb9..9abd5cb4d5c62d 100644
--- a/libc/src/math/gpu/vendor/sinf.cpp
+++ b/libc/src/math/nvptx/sinf.cpp
@@ -9,10 +9,10 @@
#include "src/math/sinf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, sinf, (float x)) { return internal::sinf(x); }
+LLVM_LIBC_FUNCTION(float, sinf, (float x)) { return __nv_sinf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/sinh.cpp b/libc/src/math/nvptx/sinh.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/sinh.cpp
rename to libc/src/math/nvptx/sinh.cpp
index be6b3ae2e2fa0b..dc6a1e16c6341e 100644
--- a/libc/src/math/gpu/vendor/sinh.cpp
+++ b/libc/src/math/nvptx/sinh.cpp
@@ -9,10 +9,10 @@
#include "src/math/sinh.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, sinh, (double x)) { return internal::sinh(x); }
+LLVM_LIBC_FUNCTION(double, sinh, (double x)) { return __nv_sinh(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/sinhf.cpp b/libc/src/math/nvptx/sinhf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/sinhf.cpp
rename to libc/src/math/nvptx/sinhf.cpp
index 99c399b62b7f6d..c9ab470ed823ce 100644
--- a/libc/src/math/gpu/vendor/sinhf.cpp
+++ b/libc/src/math/nvptx/sinhf.cpp
@@ -9,10 +9,10 @@
#include "src/math/sinhf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, sinhf, (float x)) { return internal::sinhf(x); }
+LLVM_LIBC_FUNCTION(float, sinhf, (float x)) { return __nv_sinhf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/sqrt.cpp b/libc/src/math/nvptx/sqrt.cpp
new file mode 100644
index 00000000000000..60ca5af4987b6c
--- /dev/null
+++ b/libc/src/math/nvptx/sqrt.cpp
@@ -0,0 +1,16 @@
+//===-- Implementation of the GPU sqrt function ---------------------------===//
+//
+// 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/sqrt.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, sqrt, (double x)) { return __builtin_sqrt(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/sqrtf.cpp b/libc/src/math/nvptx/sqrtf.cpp
new file mode 100644
index 00000000000000..e17f942a4d5fcf
--- /dev/null
+++ b/libc/src/math/nvptx/sqrtf.cpp
@@ -0,0 +1,16 @@
+//===-- Implementation of the GPU sqrtf function --------------------------===//
+//
+// 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/sqrtf.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, sqrtf, (float x)) { return __builtin_sqrtf(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/tan.cpp b/libc/src/math/nvptx/tan.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/tan.cpp
rename to libc/src/math/nvptx/tan.cpp
index 9a1bd9c89fcf0a..deb03dca250a51 100644
--- a/libc/src/math/gpu/vendor/tan.cpp
+++ b/libc/src/math/nvptx/tan.cpp
@@ -9,10 +9,10 @@
#include "src/math/tan.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, tan, (double x)) { return internal::tan(x); }
+LLVM_LIBC_FUNCTION(double, tan, (double x)) { return __nv_tan(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/tanf.cpp b/libc/src/math/nvptx/tanf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/tanf.cpp
rename to libc/src/math/nvptx/tanf.cpp
index a5266a8c154c5b..5739e4a1624def 100644
--- a/libc/src/math/gpu/vendor/tanf.cpp
+++ b/libc/src/math/nvptx/tanf.cpp
@@ -9,10 +9,10 @@
#include "src/math/tanf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, tanf, (float x)) { return internal::tanf(x); }
+LLVM_LIBC_FUNCTION(float, tanf, (float x)) { return __nv_tanf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/tanh.cpp b/libc/src/math/nvptx/tanh.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/tanh.cpp
rename to libc/src/math/nvptx/tanh.cpp
index 57d764f1f2a0dd..eabee2cbaf068b 100644
--- a/libc/src/math/gpu/vendor/tanh.cpp
+++ b/libc/src/math/nvptx/tanh.cpp
@@ -9,10 +9,10 @@
#include "src/math/tanh.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, tanh, (double x)) { return internal::tanh(x); }
+LLVM_LIBC_FUNCTION(double, tanh, (double x)) { return __nv_tanh(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/tanhf.cpp b/libc/src/math/nvptx/tanhf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/tanhf.cpp
rename to libc/src/math/nvptx/tanhf.cpp
index 1c9c2f3843a953..582424cb9490ae 100644
--- a/libc/src/math/gpu/vendor/tanhf.cpp
+++ b/libc/src/math/nvptx/tanhf.cpp
@@ -9,10 +9,10 @@
#include "src/math/tanhf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, tanhf, (float x)) { return internal::tanhf(x); }
+LLVM_LIBC_FUNCTION(float, tanhf, (float x)) { return __nv_tanhf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/tgamma.cpp b/libc/src/math/nvptx/tgamma.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/tgamma.cpp
rename to libc/src/math/nvptx/tgamma.cpp
index e86116a2b0abea..f92193831f9bbb 100644
--- a/libc/src/math/gpu/vendor/tgamma.cpp
+++ b/libc/src/math/nvptx/tgamma.cpp
@@ -9,10 +9,10 @@
#include "src/math/tgamma.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(double, tgamma, (double x)) { return internal::tgamma(x); }
+LLVM_LIBC_FUNCTION(double, tgamma, (double x)) { return __nv_tgamma(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/gpu/vendor/tgammaf.cpp b/libc/src/math/nvptx/tgammaf.cpp
similarity index 83%
rename from libc/src/math/gpu/vendor/tgammaf.cpp
rename to libc/src/math/nvptx/tgammaf.cpp
index 552919bae44695..833994455d5755 100644
--- a/libc/src/math/gpu/vendor/tgammaf.cpp
+++ b/libc/src/math/nvptx/tgammaf.cpp
@@ -9,10 +9,10 @@
#include "src/math/tgammaf.h"
#include "src/__support/common.h"
-#include "common.h"
+#include "declarations.h"
namespace LIBC_NAMESPACE {
-LLVM_LIBC_FUNCTION(float, tgammaf, (float x)) { return internal::tgammaf(x); }
+LLVM_LIBC_FUNCTION(float, tgammaf, (float x)) { return __nv_tgammaf(x); }
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/trunc.cpp b/libc/src/math/nvptx/trunc.cpp
new file mode 100644
index 00000000000000..773600f0f25011
--- /dev/null
+++ b/libc/src/math/nvptx/trunc.cpp
@@ -0,0 +1,16 @@
+//===-- Implementation of the GPU trunc function --------------------------===//
+//
+// 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/trunc.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(double, trunc, (double x)) { return __builtin_trunc(x); }
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/math/nvptx/truncf.cpp b/libc/src/math/nvptx/truncf.cpp
new file mode 100644
index 00000000000000..534797a3e5860c
--- /dev/null
+++ b/libc/src/math/nvptx/truncf.cpp
@@ -0,0 +1,16 @@
+//===-- Implementation of the GPU truncf function -------------------------===//
+//
+// 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/truncf.h"
+#include "src/__support/common.h"
+
+namespace LIBC_NAMESPACE {
+
+LLVM_LIBC_FUNCTION(float, truncf, (float x)) { return __builtin_truncf(x); }
+
+} // namespace LIBC_NAMESPACE
More information about the libc-commits
mailing list