[clang] [libc] [lld] [llvm] Ugly example of building libc to spirv (PR #128585)
Jon Chesterfield via cfe-commits
cfe-commits at lists.llvm.org
Mon Feb 24 13:55:56 PST 2025
https://github.com/JonChesterfield created https://github.com/llvm/llvm-project/pull/128585
Hopefully we can do better before landing. Posting for discussion.
Didn't learn a huge amount. spirv rejects the clang::amdgpu_kernel annotation and I can't work out how that's supposed to be spelled.
Libc doesn't do anything with address spaces which upsets spirv (unlike openmp which does lots)
The libc build streams bitcode through lld which refuses to do anything with it, but that can be patched.
Port is mostly copying the amdgpu folders to spirv64. Not happy with cmake assuming that spirv64 means amdgpu but also don't understand how to stop it making that assumption.
Lots of libm seems to be dead. I'm hoping Joseph will delete said dead code as it makes this patch look excessively large.
Also doesn't seem ideal to have copy&pasted the config entrypoints stuff.
Don't have a loader, probably reasonable to write one that calls llvm to jit spirv though.
>From a8970f60a98086041e3ea3c02573c48a0722e3b5 Mon Sep 17 00:00:00 2001
From: Jon Chesterfield <jonathanchesterfield at gmail.com>
Date: Mon, 24 Feb 2025 19:35:23 +0000
Subject: [PATCH 1/2] partway through porting, implicit address space
conversion is used all over devicertl
---
offload/DeviceRTL/CMakeLists.txt | 5 +++++
offload/DeviceRTL/include/State.h | 27 ++++++++++++++++++---------
offload/DeviceRTL/src/State.cpp | 21 +++++++++++++--------
offload/cmake/caches/Offload.cmake | 10 +++++-----
offload/test/offloading/bug64959.c | 2 +-
5 files changed, 42 insertions(+), 23 deletions(-)
diff --git a/offload/DeviceRTL/CMakeLists.txt b/offload/DeviceRTL/CMakeLists.txt
index 8f2a1fd01fabc..db38b017fb7e8 100644
--- a/offload/DeviceRTL/CMakeLists.txt
+++ b/offload/DeviceRTL/CMakeLists.txt
@@ -98,6 +98,7 @@ list(TRANSFORM LIBOMPTARGET_LLVM_INCLUDE_DIRS_DEVICERTL PREPEND "-I")
set(bc_flags -c -flto -std=c++17 -fvisibility=hidden
${clang_opt_flags} -nogpulib -nostdlibinc
-fno-rtti -fno-exceptions -fconvergent-functions
+ -Wno-atomic-alignment # spir-v complains about 4 byte atomics then Werror :(
-Wno-unknown-cuda-version
-DOMPTARGET_DEVICE_RUNTIME
-I${include_directory}
@@ -257,6 +258,10 @@ endfunction()
add_custom_target(omptarget.devicertl.amdgpu)
compileDeviceRTLLibrary(amdgpu amdgcn-amd-amdhsa -Xclang -mcode-object-version=none)
+add_custom_target(omptarget.devicertl.amdgpu-spirv)
+compileDeviceRTLLibrary(amdgpu-spirv spirv64-amd-amdhsa -Xclang -mcode-object-version=none)
+
+
add_custom_target(omptarget.devicertl.nvptx)
compileDeviceRTLLibrary(nvptx nvptx64-nvidia-cuda --cuda-feature=+ptx63)
diff --git a/offload/DeviceRTL/include/State.h b/offload/DeviceRTL/include/State.h
index db396dae6e445..54bcb9cf79d6b 100644
--- a/offload/DeviceRTL/include/State.h
+++ b/offload/DeviceRTL/include/State.h
@@ -102,12 +102,21 @@ struct ThreadStateTy {
ThreadStateTy *PreviousThreadState;
void init() {
- ICVState = TeamState.ICVState;
+ // assignment relies on implicit conversion between address spaces
+ // ICVState = TeamState.ICVState;
+ __builtin_memcpy(&ICVState,
+ &TeamState.ICVState,
+ sizeof(ICVState));
PreviousThreadState = nullptr;
}
void init(ThreadStateTy *PreviousTS) {
- ICVState = PreviousTS ? PreviousTS->ICVState : TeamState.ICVState;
+ __builtin_memcpy(&ICVState,
+ PreviousTS ?
+ (state::ICVStateTy*) & PreviousTS->ICVState
+ : (state::ICVStateTy*) &TeamState.ICVState,
+ sizeof(ICVState));
+
PreviousThreadState = PreviousTS;
}
};
@@ -159,7 +168,7 @@ void resetStateForThread(uint32_t TId);
{ \
if (OMP_LIKELY(ForceTeamState || !config::mayUseThreadStates() || \
!TeamState.HasThreadState)) \
- return TeamState.ICVState.Member; \
+ return (uint32_t &)TeamState.ICVState.Member; \
uint32_t TId = mapping::getThreadIdInBlock(); \
if (OMP_UNLIKELY(!ThreadStates[TId])) { \
ThreadStates[TId] = reinterpret_cast<ThreadStateTy *>( \
@@ -169,7 +178,7 @@ void resetStateForThread(uint32_t TId);
TeamState.HasThreadState = true; \
ThreadStates[TId]->init(); \
} \
- return ThreadStates[TId]->ICVState.Member; \
+ return (uint32_t &)ThreadStates[TId]->ICVState.Member; \
}
// FIXME: https://github.com/llvm/llvm-project/issues/123241.
@@ -178,8 +187,8 @@ void resetStateForThread(uint32_t TId);
auto TId = mapping::getThreadIdInBlock(); \
if (OMP_UNLIKELY(!ForceTeamState && config::mayUseThreadStates() && \
TeamState.HasThreadState && ThreadStates[TId])) \
- return ThreadStates[TId]->ICVState.Member; \
- return TeamState.ICVState.Member; \
+ return (uint32_t &)ThreadStates[TId]->ICVState.Member; \
+ return (uint32_t &)TeamState.ICVState.Member; \
}
[[gnu::always_inline, gnu::flatten]] inline uint32_t &
@@ -210,9 +219,9 @@ lookup32(ValueKind Kind, bool IsReadonly, IdentTy *Ident, bool ForceTeamState) {
lookupImpl(RunSchedChunkVar, ForceTeamState);
lookupForModify32Impl(RunSchedChunkVar, Ident, ForceTeamState);
case state::VK_ParallelTeamSize:
- return TeamState.ParallelTeamSize;
+ return (uint32_t &)TeamState.ParallelTeamSize;
case state::VK_HasThreadState:
- return TeamState.HasThreadState;
+ return (uint32_t &)TeamState.HasThreadState;
default:
break;
}
@@ -223,7 +232,7 @@ lookup32(ValueKind Kind, bool IsReadonly, IdentTy *Ident, bool ForceTeamState) {
lookupPtr(ValueKind Kind, bool IsReadonly, bool ForceTeamState) {
switch (Kind) {
case state::VK_ParallelRegionFn:
- return TeamState.ParallelRegionFnVar;
+ return (void *&)TeamState.ParallelRegionFnVar;
default:
break;
}
diff --git a/offload/DeviceRTL/src/State.cpp b/offload/DeviceRTL/src/State.cpp
index cbe9735145340..fe7fab0ef8edb 100644
--- a/offload/DeviceRTL/src/State.cpp
+++ b/offload/DeviceRTL/src/State.cpp
@@ -160,7 +160,7 @@ void SharedMemorySmartStackTy::pop(void *Ptr, uint64_t Bytes) {
} // namespace
-void *memory::getDynamicBuffer() { return DynamicSharedBuffer; }
+void *memory::getDynamicBuffer() { return (void *)DynamicSharedBuffer; }
void *memory::allocShared(uint64_t Bytes, const char *Reason) {
return SharedMemorySmartStack.push(Bytes);
@@ -219,6 +219,7 @@ bool state::TeamStateTy::operator==(const TeamStateTy &Other) const {
void state::TeamStateTy::assertEqual(TeamStateTy &Other) const {
ICVState.assertEqual(Other.ICVState);
+ // Other.ICVState.assertEqual(ICVState);
ASSERT(ParallelTeamSize == Other.ParallelTeamSize, nullptr);
ASSERT(HasThreadState == Other.HasThreadState, nullptr);
}
@@ -273,7 +274,7 @@ void state::enterDataEnvironment(IdentTy *Ident) {
unsigned TId = mapping::getThreadIdInBlock();
ThreadStateTy *NewThreadState = static_cast<ThreadStateTy *>(
memory::allocGlobal(sizeof(ThreadStateTy), "ThreadStates alloc"));
- uintptr_t *ThreadStatesBitsPtr = reinterpret_cast<uintptr_t *>(&ThreadStates);
+ uintptr_t *ThreadStatesBitsPtr = (uintptr_t *)(&ThreadStates);
if (!atomic::load(ThreadStatesBitsPtr, atomic::seq_cst)) {
uint32_t Bytes =
sizeof(ThreadStates[0]) * mapping::getNumberOfThreadsInBlock();
@@ -313,18 +314,22 @@ void state::resetStateForThread(uint32_t TId) {
}
void state::runAndCheckState(void(Func(void))) {
- TeamStateTy OldTeamState = TeamState;
- OldTeamState.assertEqual(TeamState);
+ // TeamStateTy OldTeamState = TeamState;
+ TeamStateTy OldTeamState;
+ __builtin_memcpy(&OldTeamState,
+ &TeamState,
+ sizeof(TeamStateTy));
+ OldTeamState.assertEqual((TeamStateTy&)TeamState);
Func();
- OldTeamState.assertEqual(TeamState);
+ OldTeamState.assertEqual((TeamStateTy&)TeamState);
}
void state::assumeInitialState(bool IsSPMD) {
TeamStateTy InitialTeamState;
InitialTeamState.init(IsSPMD);
- InitialTeamState.assertEqual(TeamState);
+ InitialTeamState.assertEqual((TeamStateTy&)TeamState);
ASSERT(mapping::isSPMDMode() == IsSPMD, nullptr);
}
@@ -461,7 +466,7 @@ constexpr uint64_t NUM_SHARED_VARIABLES_IN_SHARED_MEM = 64;
void __kmpc_begin_sharing_variables(void ***GlobalArgs, uint64_t nArgs) {
if (nArgs <= NUM_SHARED_VARIABLES_IN_SHARED_MEM) {
- SharedMemVariableSharingSpacePtr = &SharedMemVariableSharingSpace[0];
+ SharedMemVariableSharingSpacePtr = (void**)&SharedMemVariableSharingSpace[0];
} else {
SharedMemVariableSharingSpacePtr = (void **)memory::allocGlobal(
nArgs * sizeof(void *), "new extended args");
@@ -472,7 +477,7 @@ void __kmpc_begin_sharing_variables(void ***GlobalArgs, uint64_t nArgs) {
}
void __kmpc_end_sharing_variables() {
- if (SharedMemVariableSharingSpacePtr != &SharedMemVariableSharingSpace[0])
+ if ((void*)SharedMemVariableSharingSpacePtr != (void*)&SharedMemVariableSharingSpace[0])
memory::freeGlobal(SharedMemVariableSharingSpacePtr, "new extended args");
}
diff --git a/offload/cmake/caches/Offload.cmake b/offload/cmake/caches/Offload.cmake
index 5533a6508f5d5..6c74fc1d4b762 100644
--- a/offload/cmake/caches/Offload.cmake
+++ b/offload/cmake/caches/Offload.cmake
@@ -1,9 +1,9 @@
set(LLVM_ENABLE_PROJECTS "clang;clang-tools-extra;lld" CACHE STRING "")
-set(LLVM_ENABLE_RUNTIMES "compiler-rt;libunwind;libcxx;libcxxabi;openmp;offload" CACHE STRING "")
+set(LLVM_ENABLE_RUNTIMES "compiler-rt;openmp;offload" CACHE STRING "")
set(LLVM_ENABLE_PER_TARGET_RUNTIME_DIR ON CACHE BOOL "")
set(LLVM_RUNTIME_TARGETS default;amdgcn-amd-amdhsa;nvptx64-nvidia-cuda CACHE STRING "")
-set(RUNTIMES_nvptx64-nvidia-cuda_CACHE_FILES "${CMAKE_SOURCE_DIR}/../libcxx/cmake/caches/NVPTX.cmake" CACHE STRING "")
-set(RUNTIMES_amdgcn-amd-amdhsa_CACHE_FILES "${CMAKE_SOURCE_DIR}/../libcxx/cmake/caches/AMDGPU.cmake" CACHE STRING "")
-set(RUNTIMES_nvptx64-nvidia-cuda_LLVM_ENABLE_RUNTIMES "compiler-rt;libc;libcxx;libcxxabi" CACHE STRING "")
-set(RUNTIMES_amdgcn-amd-amdhsa_LLVM_ENABLE_RUNTIMES "compiler-rt;libc;libcxx;libcxxabi" CACHE STRING "")
+# set(RUNTIMES_nvptx64-nvidia-cuda_CACHE_FILES "${CMAKE_SOURCE_DIR}/../libcxx/cmake/caches/NVPTX.cmake" CACHE STRING "")
+# set(RUNTIMES_amdgcn-amd-amdhsa_CACHE_FILES "${CMAKE_SOURCE_DIR}/../libcxx/cmake/caches/AMDGPU.cmake" CACHE STRING "")
+set(RUNTIMES_nvptx64-nvidia-cuda_LLVM_ENABLE_RUNTIMES "compiler-rt;libc" CACHE STRING "")
+set(RUNTIMES_amdgcn-amd-amdhsa_LLVM_ENABLE_RUNTIMES "compiler-rt;libc" CACHE STRING "")
diff --git a/offload/test/offloading/bug64959.c b/offload/test/offloading/bug64959.c
index eddc55325ffe9..89d523bdf6de9 100644
--- a/offload/test/offloading/bug64959.c
+++ b/offload/test/offloading/bug64959.c
@@ -55,6 +55,6 @@ int main(void) {
return 1;
}
// CHECK: Success
- printf("Success\n");
+ printf("Suc disabled cess\n");
return 0;
}
>From 874f408032cdb12188810a06c9b68d90c35deecb Mon Sep 17 00:00:00 2001
From: Jon Chesterfield <jonathanchesterfield at gmail.com>
Date: Mon, 24 Feb 2025 21:52:52 +0000
Subject: [PATCH 2/2] port libc instead, nothing much discovered
---
clang/lib/Basic/Targets/SPIR.h | 2 +
.../cmake/modules/LLVMLibCArchitectures.cmake | 4 +
libc/config/gpu/spirv64/config.json | 37 ++
libc/config/gpu/spirv64/entrypoints.txt | 607 ++++++++++++++++++
libc/config/gpu/spirv64/headers.txt | 21 +
libc/src/math/spirv64/CMakeLists.txt | 541 ++++++++++++++++
libc/src/math/spirv64/acos.cpp | 19 +
libc/src/math/spirv64/acosf.cpp | 19 +
libc/src/math/spirv64/acosh.cpp | 19 +
libc/src/math/spirv64/acoshf.cpp | 19 +
libc/src/math/spirv64/asin.cpp | 19 +
libc/src/math/spirv64/asinf.cpp | 19 +
libc/src/math/spirv64/asinh.cpp | 19 +
libc/src/math/spirv64/asinhf.cpp | 19 +
libc/src/math/spirv64/atan.cpp | 19 +
libc/src/math/spirv64/atan2.cpp | 21 +
libc/src/math/spirv64/atan2f.cpp | 21 +
libc/src/math/spirv64/atanf.cpp | 19 +
libc/src/math/spirv64/atanh.cpp | 19 +
libc/src/math/spirv64/atanhf.cpp | 19 +
libc/src/math/spirv64/ceil.cpp | 17 +
libc/src/math/spirv64/ceilf.cpp | 17 +
libc/src/math/spirv64/copysign.cpp | 19 +
libc/src/math/spirv64/copysignf.cpp | 19 +
libc/src/math/spirv64/cos.cpp | 19 +
libc/src/math/spirv64/cosf.cpp | 19 +
libc/src/math/spirv64/cosh.cpp | 19 +
libc/src/math/spirv64/coshf.cpp | 19 +
libc/src/math/spirv64/declarations.h | 91 +++
libc/src/math/spirv64/erf.cpp | 19 +
libc/src/math/spirv64/erff.cpp | 19 +
libc/src/math/spirv64/exp.cpp | 19 +
libc/src/math/spirv64/exp10.cpp | 19 +
libc/src/math/spirv64/exp10f.cpp | 19 +
libc/src/math/spirv64/exp2.cpp | 19 +
libc/src/math/spirv64/exp2f.cpp | 19 +
libc/src/math/spirv64/expf.cpp | 19 +
libc/src/math/spirv64/expm1.cpp | 19 +
libc/src/math/spirv64/expm1f.cpp | 19 +
libc/src/math/spirv64/fabs.cpp | 17 +
libc/src/math/spirv64/fabsf.cpp | 17 +
libc/src/math/spirv64/fdim.cpp | 21 +
libc/src/math/spirv64/fdimf.cpp | 21 +
libc/src/math/spirv64/floor.cpp | 17 +
libc/src/math/spirv64/floorf.cpp | 17 +
libc/src/math/spirv64/fma.cpp | 19 +
libc/src/math/spirv64/fmaf.cpp | 19 +
libc/src/math/spirv64/fmax.cpp | 22 +
libc/src/math/spirv64/fmaxf.cpp | 20 +
libc/src/math/spirv64/fmin.cpp | 20 +
libc/src/math/spirv64/fminf.cpp | 20 +
libc/src/math/spirv64/fmod.cpp | 19 +
libc/src/math/spirv64/fmodf.cpp | 19 +
libc/src/math/spirv64/frexp.cpp | 21 +
libc/src/math/spirv64/frexpf.cpp | 21 +
libc/src/math/spirv64/hypot.cpp | 21 +
libc/src/math/spirv64/hypotf.cpp | 21 +
libc/src/math/spirv64/ilogb.cpp | 19 +
libc/src/math/spirv64/ilogbf.cpp | 19 +
libc/src/math/spirv64/ldexp.cpp | 21 +
libc/src/math/spirv64/ldexpf.cpp | 21 +
libc/src/math/spirv64/lgamma.cpp | 19 +
libc/src/math/spirv64/lgamma_r.cpp | 24 +
libc/src/math/spirv64/llrint.cpp | 21 +
libc/src/math/spirv64/llrintf.cpp | 21 +
libc/src/math/spirv64/log.cpp | 19 +
libc/src/math/spirv64/log10.cpp | 19 +
libc/src/math/spirv64/log10f.cpp | 19 +
libc/src/math/spirv64/log1p.cpp | 19 +
libc/src/math/spirv64/log1pf.cpp | 19 +
libc/src/math/spirv64/log2.cpp | 19 +
libc/src/math/spirv64/log2f.cpp | 19 +
libc/src/math/spirv64/logb.cpp | 19 +
libc/src/math/spirv64/logbf.cpp | 19 +
libc/src/math/spirv64/logf.cpp | 19 +
libc/src/math/spirv64/lrint.cpp | 21 +
libc/src/math/spirv64/lrintf.cpp | 21 +
libc/src/math/spirv64/nearbyint.cpp | 19 +
libc/src/math/spirv64/nearbyintf.cpp | 19 +
libc/src/math/spirv64/nextafter.cpp | 21 +
libc/src/math/spirv64/nextafterf.cpp | 21 +
libc/src/math/spirv64/platform.h | 55 ++
libc/src/math/spirv64/powf.cpp | 21 +
libc/src/math/spirv64/powi.cpp | 21 +
libc/src/math/spirv64/powif.cpp | 21 +
libc/src/math/spirv64/remainder.cpp | 19 +
libc/src/math/spirv64/remainderf.cpp | 19 +
libc/src/math/spirv64/remquo.cpp | 24 +
libc/src/math/spirv64/remquof.cpp | 24 +
libc/src/math/spirv64/rint.cpp | 17 +
libc/src/math/spirv64/rintf.cpp | 17 +
libc/src/math/spirv64/round.cpp | 17 +
libc/src/math/spirv64/roundf.cpp | 17 +
libc/src/math/spirv64/scalbn.cpp | 21 +
libc/src/math/spirv64/scalbnf.cpp | 21 +
libc/src/math/spirv64/sin.cpp | 19 +
libc/src/math/spirv64/sincos.cpp | 21 +
libc/src/math/spirv64/sincosf.cpp | 21 +
libc/src/math/spirv64/sinf.cpp | 19 +
libc/src/math/spirv64/sinh.cpp | 19 +
libc/src/math/spirv64/sinhf.cpp | 19 +
libc/src/math/spirv64/sqrt.cpp | 17 +
libc/src/math/spirv64/sqrtf.cpp | 17 +
libc/src/math/spirv64/tan.cpp | 19 +
libc/src/math/spirv64/tanf.cpp | 19 +
libc/src/math/spirv64/tanh.cpp | 19 +
libc/src/math/spirv64/tanhf.cpp | 19 +
libc/src/math/spirv64/tgamma.cpp | 19 +
libc/src/math/spirv64/tgammaf.cpp | 19 +
libc/src/math/spirv64/trunc.cpp | 17 +
libc/src/math/spirv64/truncf.cpp | 17 +
libc/startup/gpu/spirv64/CMakeLists.txt | 15 +
libc/startup/gpu/spirv64/start.cpp | 80 +++
lld/ELF/InputFiles.cpp | 6 +-
offload/cmake/caches/Offload.cmake | 4 +-
115 files changed, 3457 insertions(+), 2 deletions(-)
create mode 100644 libc/config/gpu/spirv64/config.json
create mode 100644 libc/config/gpu/spirv64/entrypoints.txt
create mode 100644 libc/config/gpu/spirv64/headers.txt
create mode 100644 libc/src/math/spirv64/CMakeLists.txt
create mode 100644 libc/src/math/spirv64/acos.cpp
create mode 100644 libc/src/math/spirv64/acosf.cpp
create mode 100644 libc/src/math/spirv64/acosh.cpp
create mode 100644 libc/src/math/spirv64/acoshf.cpp
create mode 100644 libc/src/math/spirv64/asin.cpp
create mode 100644 libc/src/math/spirv64/asinf.cpp
create mode 100644 libc/src/math/spirv64/asinh.cpp
create mode 100644 libc/src/math/spirv64/asinhf.cpp
create mode 100644 libc/src/math/spirv64/atan.cpp
create mode 100644 libc/src/math/spirv64/atan2.cpp
create mode 100644 libc/src/math/spirv64/atan2f.cpp
create mode 100644 libc/src/math/spirv64/atanf.cpp
create mode 100644 libc/src/math/spirv64/atanh.cpp
create mode 100644 libc/src/math/spirv64/atanhf.cpp
create mode 100644 libc/src/math/spirv64/ceil.cpp
create mode 100644 libc/src/math/spirv64/ceilf.cpp
create mode 100644 libc/src/math/spirv64/copysign.cpp
create mode 100644 libc/src/math/spirv64/copysignf.cpp
create mode 100644 libc/src/math/spirv64/cos.cpp
create mode 100644 libc/src/math/spirv64/cosf.cpp
create mode 100644 libc/src/math/spirv64/cosh.cpp
create mode 100644 libc/src/math/spirv64/coshf.cpp
create mode 100644 libc/src/math/spirv64/declarations.h
create mode 100644 libc/src/math/spirv64/erf.cpp
create mode 100644 libc/src/math/spirv64/erff.cpp
create mode 100644 libc/src/math/spirv64/exp.cpp
create mode 100644 libc/src/math/spirv64/exp10.cpp
create mode 100644 libc/src/math/spirv64/exp10f.cpp
create mode 100644 libc/src/math/spirv64/exp2.cpp
create mode 100644 libc/src/math/spirv64/exp2f.cpp
create mode 100644 libc/src/math/spirv64/expf.cpp
create mode 100644 libc/src/math/spirv64/expm1.cpp
create mode 100644 libc/src/math/spirv64/expm1f.cpp
create mode 100644 libc/src/math/spirv64/fabs.cpp
create mode 100644 libc/src/math/spirv64/fabsf.cpp
create mode 100644 libc/src/math/spirv64/fdim.cpp
create mode 100644 libc/src/math/spirv64/fdimf.cpp
create mode 100644 libc/src/math/spirv64/floor.cpp
create mode 100644 libc/src/math/spirv64/floorf.cpp
create mode 100644 libc/src/math/spirv64/fma.cpp
create mode 100644 libc/src/math/spirv64/fmaf.cpp
create mode 100644 libc/src/math/spirv64/fmax.cpp
create mode 100644 libc/src/math/spirv64/fmaxf.cpp
create mode 100644 libc/src/math/spirv64/fmin.cpp
create mode 100644 libc/src/math/spirv64/fminf.cpp
create mode 100644 libc/src/math/spirv64/fmod.cpp
create mode 100644 libc/src/math/spirv64/fmodf.cpp
create mode 100644 libc/src/math/spirv64/frexp.cpp
create mode 100644 libc/src/math/spirv64/frexpf.cpp
create mode 100644 libc/src/math/spirv64/hypot.cpp
create mode 100644 libc/src/math/spirv64/hypotf.cpp
create mode 100644 libc/src/math/spirv64/ilogb.cpp
create mode 100644 libc/src/math/spirv64/ilogbf.cpp
create mode 100644 libc/src/math/spirv64/ldexp.cpp
create mode 100644 libc/src/math/spirv64/ldexpf.cpp
create mode 100644 libc/src/math/spirv64/lgamma.cpp
create mode 100644 libc/src/math/spirv64/lgamma_r.cpp
create mode 100644 libc/src/math/spirv64/llrint.cpp
create mode 100644 libc/src/math/spirv64/llrintf.cpp
create mode 100644 libc/src/math/spirv64/log.cpp
create mode 100644 libc/src/math/spirv64/log10.cpp
create mode 100644 libc/src/math/spirv64/log10f.cpp
create mode 100644 libc/src/math/spirv64/log1p.cpp
create mode 100644 libc/src/math/spirv64/log1pf.cpp
create mode 100644 libc/src/math/spirv64/log2.cpp
create mode 100644 libc/src/math/spirv64/log2f.cpp
create mode 100644 libc/src/math/spirv64/logb.cpp
create mode 100644 libc/src/math/spirv64/logbf.cpp
create mode 100644 libc/src/math/spirv64/logf.cpp
create mode 100644 libc/src/math/spirv64/lrint.cpp
create mode 100644 libc/src/math/spirv64/lrintf.cpp
create mode 100644 libc/src/math/spirv64/nearbyint.cpp
create mode 100644 libc/src/math/spirv64/nearbyintf.cpp
create mode 100644 libc/src/math/spirv64/nextafter.cpp
create mode 100644 libc/src/math/spirv64/nextafterf.cpp
create mode 100644 libc/src/math/spirv64/platform.h
create mode 100644 libc/src/math/spirv64/powf.cpp
create mode 100644 libc/src/math/spirv64/powi.cpp
create mode 100644 libc/src/math/spirv64/powif.cpp
create mode 100644 libc/src/math/spirv64/remainder.cpp
create mode 100644 libc/src/math/spirv64/remainderf.cpp
create mode 100644 libc/src/math/spirv64/remquo.cpp
create mode 100644 libc/src/math/spirv64/remquof.cpp
create mode 100644 libc/src/math/spirv64/rint.cpp
create mode 100644 libc/src/math/spirv64/rintf.cpp
create mode 100644 libc/src/math/spirv64/round.cpp
create mode 100644 libc/src/math/spirv64/roundf.cpp
create mode 100644 libc/src/math/spirv64/scalbn.cpp
create mode 100644 libc/src/math/spirv64/scalbnf.cpp
create mode 100644 libc/src/math/spirv64/sin.cpp
create mode 100644 libc/src/math/spirv64/sincos.cpp
create mode 100644 libc/src/math/spirv64/sincosf.cpp
create mode 100644 libc/src/math/spirv64/sinf.cpp
create mode 100644 libc/src/math/spirv64/sinh.cpp
create mode 100644 libc/src/math/spirv64/sinhf.cpp
create mode 100644 libc/src/math/spirv64/sqrt.cpp
create mode 100644 libc/src/math/spirv64/sqrtf.cpp
create mode 100644 libc/src/math/spirv64/tan.cpp
create mode 100644 libc/src/math/spirv64/tanf.cpp
create mode 100644 libc/src/math/spirv64/tanh.cpp
create mode 100644 libc/src/math/spirv64/tanhf.cpp
create mode 100644 libc/src/math/spirv64/tgamma.cpp
create mode 100644 libc/src/math/spirv64/tgammaf.cpp
create mode 100644 libc/src/math/spirv64/trunc.cpp
create mode 100644 libc/src/math/spirv64/truncf.cpp
create mode 100644 libc/startup/gpu/spirv64/CMakeLists.txt
create mode 100644 libc/startup/gpu/spirv64/start.cpp
diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h
index 61f9ef7e3e361..610efa1fe00d9 100644
--- a/clang/lib/Basic/Targets/SPIR.h
+++ b/clang/lib/Basic/Targets/SPIR.h
@@ -399,6 +399,8 @@ class LLVM_LIBRARY_VISIBILITY SPIRV64AMDGCNTargetInfo final
HasLegalHalfType = true;
HasFloat16 = true;
HalfArgsAndReturns = true;
+
+ MaxAtomicPromoteWidth = MaxAtomicInlineWidth = 64;
}
bool hasBFloat16Type() const override { return true; }
diff --git a/libc/cmake/modules/LLVMLibCArchitectures.cmake b/libc/cmake/modules/LLVMLibCArchitectures.cmake
index fbb1091ddabab..8d74f71d5cba0 100644
--- a/libc/cmake/modules/LLVMLibCArchitectures.cmake
+++ b/libc/cmake/modules/LLVMLibCArchitectures.cmake
@@ -49,6 +49,8 @@ function(get_arch_and_system_from_triple triple arch_var sys_var)
set(target_arch "riscv32")
elseif(target_arch MATCHES "^riscv64")
set(target_arch "riscv64")
+ elseif(target_arch MATCHES "^spirv64")
+ set(target_arch "spirv64")
elseif(target_arch MATCHES "^amdgcn")
set(target_arch "amdgpu")
elseif(target_arch MATCHES "^nvptx64")
@@ -160,6 +162,8 @@ elseif(LIBC_TARGET_ARCHITECTURE STREQUAL "riscv32")
set(LIBC_TARGET_ARCHITECTURE "riscv")
elseif(LIBC_TARGET_ARCHITECTURE STREQUAL "amdgpu")
set(LIBC_TARGET_ARCHITECTURE_IS_AMDGPU TRUE)
+elseif(LIBC_TARGET_ARCHITECTURE STREQUAL "spirv64")
+ set(LIBC_TARGET_ARCHITECTURE_IS_SPIRV64 TRUE)
elseif(LIBC_TARGET_ARCHITECTURE STREQUAL "nvptx")
set(LIBC_TARGET_ARCHITECTURE_IS_NVPTX TRUE)
else()
diff --git a/libc/config/gpu/spirv64/config.json b/libc/config/gpu/spirv64/config.json
new file mode 100644
index 0000000000000..d99f48ecbede1
--- /dev/null
+++ b/libc/config/gpu/spirv64/config.json
@@ -0,0 +1,37 @@
+{
+ "errno": {
+ "LIBC_CONF_ERRNO_MODE": {
+ "value": "LIBC_ERRNO_MODE_SHARED"
+ }
+ },
+ "printf": {
+ "LIBC_CONF_PRINTF_DISABLE_FLOAT": {
+ "value": true
+ },
+ "LIBC_CONF_PRINTF_DISABLE_INDEX_MODE": {
+ "value": true
+ },
+ "LIBC_CONF_PRINTF_DISABLE_WRITE_INT": {
+ "value": true
+ },
+ "LIBC_CONF_PRINTF_FLOAT_TO_STR_USE_MEGA_LONG_DOUBLE_TABLE": {
+ "value": false
+ },
+ "LIBC_CONF_PRINTF_DISABLE_STRERROR": {
+ "value": true
+ }
+ },
+ "scanf": {
+ "LIBC_CONF_SCANF_DISABLE_FLOAT": {
+ "value": true
+ },
+ "LIBC_CONF_SCANF_DISABLE_INDEX_MODE": {
+ "value": true
+ }
+ },
+ "math": {
+ "LIBC_CONF_MATH_OPTIMIZATIONS": {
+ "value": "(LIBC_MATH_SKIP_ACCURATE_PASS | LIBC_MATH_SMALL_TABLES | LIBC_MATH_NO_ERRNO | LIBC_MATH_NO_EXCEPT)"
+ }
+ }
+}
diff --git a/libc/config/gpu/spirv64/entrypoints.txt b/libc/config/gpu/spirv64/entrypoints.txt
new file mode 100644
index 0000000000000..291d86b4dd587
--- /dev/null
+++ b/libc/config/gpu/spirv64/entrypoints.txt
@@ -0,0 +1,607 @@
+set(TARGET_LIBC_ENTRYPOINTS
+ # assert.h entrypoints
+ libc.src.assert.__assert_fail
+
+ # ctype.h entrypoints
+ libc.src.ctype.isalnum
+ libc.src.ctype.isalnum_l
+ libc.src.ctype.isalpha
+ libc.src.ctype.isalpha_l
+ libc.src.ctype.isascii
+ libc.src.ctype.isblank
+ libc.src.ctype.isblank_l
+ libc.src.ctype.iscntrl
+ libc.src.ctype.iscntrl_l
+ libc.src.ctype.isdigit
+ libc.src.ctype.isdigit_l
+ libc.src.ctype.isgraph
+ libc.src.ctype.isgraph_l
+ libc.src.ctype.islower
+ libc.src.ctype.islower_l
+ libc.src.ctype.isprint
+ libc.src.ctype.isprint_l
+ libc.src.ctype.ispunct
+ libc.src.ctype.ispunct_l
+ libc.src.ctype.isspace
+ libc.src.ctype.isspace_l
+ libc.src.ctype.isupper
+ libc.src.ctype.isupper_l
+ libc.src.ctype.isxdigit
+ libc.src.ctype.isxdigit_l
+ libc.src.ctype.toascii
+ libc.src.ctype.tolower
+ libc.src.ctype.tolower_l
+ libc.src.ctype.toupper
+ libc.src.ctype.toupper_l
+
+ # string.h entrypoints
+ libc.src.string.memccpy
+ libc.src.string.memchr
+ libc.src.string.memcmp
+ libc.src.string.memcpy
+ libc.src.string.memmem
+ libc.src.string.memmove
+ libc.src.string.mempcpy
+ libc.src.string.memrchr
+ libc.src.string.memset
+ libc.src.string.stpcpy
+ libc.src.string.stpncpy
+ libc.src.string.strcasestr
+ libc.src.string.strcat
+ libc.src.string.strchr
+ libc.src.string.strchrnul
+ libc.src.string.strcmp
+ libc.src.string.strcoll
+ libc.src.string.strcoll_l
+ libc.src.string.strcpy
+ libc.src.string.strcspn
+ libc.src.string.strdup
+ libc.src.string.strerror
+ libc.src.string.strlcat
+ libc.src.string.strlcpy
+ libc.src.string.strlen
+ libc.src.string.strncat
+ libc.src.string.strncmp
+ libc.src.string.strncpy
+ libc.src.string.strndup
+ libc.src.string.strnlen
+ libc.src.string.strpbrk
+ libc.src.string.strrchr
+ libc.src.string.strsep
+ libc.src.string.strspn
+ libc.src.string.strstr
+ libc.src.string.strtok
+ libc.src.string.strtok_r
+ libc.src.string.strxfrm
+ libc.src.string.strxfrm_l
+
+ # strings.h entrypoints
+ libc.src.strings.bcmp
+ libc.src.strings.bcopy
+ libc.src.strings.bzero
+ libc.src.strings.index
+ libc.src.strings.rindex
+ libc.src.strings.strcasecmp
+ libc.src.strings.strncasecmp
+
+ # stdbit.h entrypoints
+ libc.src.stdbit.stdc_bit_ceil_uc
+ libc.src.stdbit.stdc_bit_ceil_ui
+ libc.src.stdbit.stdc_bit_ceil_ul
+ libc.src.stdbit.stdc_bit_ceil_ull
+ libc.src.stdbit.stdc_bit_ceil_us
+ libc.src.stdbit.stdc_bit_floor_uc
+ libc.src.stdbit.stdc_bit_floor_ui
+ libc.src.stdbit.stdc_bit_floor_ul
+ libc.src.stdbit.stdc_bit_floor_ull
+ libc.src.stdbit.stdc_bit_floor_us
+ libc.src.stdbit.stdc_bit_width_uc
+ libc.src.stdbit.stdc_bit_width_ui
+ libc.src.stdbit.stdc_bit_width_ul
+ libc.src.stdbit.stdc_bit_width_ull
+ libc.src.stdbit.stdc_bit_width_us
+ libc.src.stdbit.stdc_count_ones_uc
+ libc.src.stdbit.stdc_count_ones_ui
+ libc.src.stdbit.stdc_count_ones_ul
+ libc.src.stdbit.stdc_count_ones_ull
+ libc.src.stdbit.stdc_count_ones_us
+ libc.src.stdbit.stdc_count_zeros_uc
+ libc.src.stdbit.stdc_count_zeros_ui
+ libc.src.stdbit.stdc_count_zeros_ul
+ libc.src.stdbit.stdc_count_zeros_ull
+ libc.src.stdbit.stdc_count_zeros_us
+ libc.src.stdbit.stdc_first_leading_one_uc
+ libc.src.stdbit.stdc_first_leading_one_ui
+ libc.src.stdbit.stdc_first_leading_one_ul
+ libc.src.stdbit.stdc_first_leading_one_ull
+ libc.src.stdbit.stdc_first_leading_one_us
+ libc.src.stdbit.stdc_first_leading_zero_uc
+ libc.src.stdbit.stdc_first_leading_zero_ui
+ libc.src.stdbit.stdc_first_leading_zero_ul
+ libc.src.stdbit.stdc_first_leading_zero_ull
+ libc.src.stdbit.stdc_first_leading_zero_us
+ libc.src.stdbit.stdc_first_trailing_one_uc
+ libc.src.stdbit.stdc_first_trailing_one_ui
+ libc.src.stdbit.stdc_first_trailing_one_ul
+ libc.src.stdbit.stdc_first_trailing_one_ull
+ libc.src.stdbit.stdc_first_trailing_one_us
+ libc.src.stdbit.stdc_first_trailing_zero_uc
+ libc.src.stdbit.stdc_first_trailing_zero_ui
+ libc.src.stdbit.stdc_first_trailing_zero_ul
+ libc.src.stdbit.stdc_first_trailing_zero_ull
+ libc.src.stdbit.stdc_first_trailing_zero_us
+ libc.src.stdbit.stdc_has_single_bit_uc
+ libc.src.stdbit.stdc_has_single_bit_ui
+ libc.src.stdbit.stdc_has_single_bit_ul
+ libc.src.stdbit.stdc_has_single_bit_ull
+ libc.src.stdbit.stdc_has_single_bit_us
+ libc.src.stdbit.stdc_leading_ones_uc
+ libc.src.stdbit.stdc_leading_ones_ui
+ libc.src.stdbit.stdc_leading_ones_ul
+ libc.src.stdbit.stdc_leading_ones_ull
+ libc.src.stdbit.stdc_leading_ones_us
+ libc.src.stdbit.stdc_leading_zeros_uc
+ libc.src.stdbit.stdc_leading_zeros_ui
+ libc.src.stdbit.stdc_leading_zeros_ul
+ libc.src.stdbit.stdc_leading_zeros_ull
+ libc.src.stdbit.stdc_leading_zeros_us
+ libc.src.stdbit.stdc_trailing_ones_uc
+ libc.src.stdbit.stdc_trailing_ones_ui
+ libc.src.stdbit.stdc_trailing_ones_ul
+ libc.src.stdbit.stdc_trailing_ones_ull
+ libc.src.stdbit.stdc_trailing_ones_us
+ libc.src.stdbit.stdc_trailing_zeros_uc
+ libc.src.stdbit.stdc_trailing_zeros_ui
+ libc.src.stdbit.stdc_trailing_zeros_ul
+ libc.src.stdbit.stdc_trailing_zeros_ull
+ libc.src.stdbit.stdc_trailing_zeros_us
+
+ # stdlib.h entrypoints
+ libc.src.stdlib._Exit
+ libc.src.stdlib.abort
+ libc.src.stdlib.abs
+ libc.src.stdlib.atexit
+ libc.src.stdlib.atof
+ libc.src.stdlib.atoi
+ libc.src.stdlib.atol
+ libc.src.stdlib.atoll
+ libc.src.stdlib.bsearch
+ libc.src.stdlib.div
+ libc.src.stdlib.exit
+ libc.src.stdlib.labs
+ libc.src.stdlib.ldiv
+ libc.src.stdlib.llabs
+ libc.src.stdlib.lldiv
+ libc.src.stdlib.qsort
+ libc.src.stdlib.qsort_r
+ libc.src.stdlib.rand
+ libc.src.stdlib.srand
+ libc.src.stdlib.strtod
+ libc.src.stdlib.strtod_l
+ libc.src.stdlib.strtof
+ libc.src.stdlib.strtof_l
+ libc.src.stdlib.strtol
+ libc.src.stdlib.strtol_l
+ libc.src.stdlib.strtold
+ libc.src.stdlib.strtold_l
+ libc.src.stdlib.strtoll
+ libc.src.stdlib.strtoll_l
+ libc.src.stdlib.strtoul
+ libc.src.stdlib.strtoul_l
+ libc.src.stdlib.strtoull
+ libc.src.stdlib.strtoull_l
+ libc.src.stdlib.at_quick_exit
+ libc.src.stdlib.quick_exit
+ libc.src.stdlib.getenv
+ libc.src.stdlib.system
+
+ # TODO: Implement these correctly
+ libc.src.stdlib.aligned_alloc
+ libc.src.stdlib.calloc
+ libc.src.stdlib.free
+ libc.src.stdlib.malloc
+ libc.src.stdlib.realloc
+
+ # errno.h entrypoints
+ libc.src.errno.errno
+
+ # stdio.h entrypoints
+ libc.src.stdio.clearerr
+ libc.src.stdio.fclose
+ libc.src.stdio.printf
+ libc.src.stdio.vprintf
+ libc.src.stdio.fprintf
+ libc.src.stdio.vfprintf
+ libc.src.stdio.snprintf
+ libc.src.stdio.sprintf
+ libc.src.stdio.vsnprintf
+ libc.src.stdio.vsprintf
+ libc.src.stdio.asprintf
+ libc.src.stdio.vasprintf
+ libc.src.stdio.scanf
+ libc.src.stdio.vscanf
+ libc.src.stdio.fscanf
+ libc.src.stdio.vfscanf
+ libc.src.stdio.sscanf
+ libc.src.stdio.vsscanf
+ libc.src.stdio.feof
+ libc.src.stdio.ferror
+ libc.src.stdio.fflush
+ libc.src.stdio.fgetc
+ libc.src.stdio.fgets
+ libc.src.stdio.fopen
+ libc.src.stdio.fputc
+ libc.src.stdio.fputs
+ libc.src.stdio.fread
+ libc.src.stdio.fseek
+ libc.src.stdio.ftell
+ libc.src.stdio.fwrite
+ libc.src.stdio.getc
+ libc.src.stdio.getchar
+ libc.src.stdio.putc
+ libc.src.stdio.putchar
+ libc.src.stdio.puts
+ libc.src.stdio.remove
+ libc.src.stdio.rename
+ libc.src.stdio.stderr
+ libc.src.stdio.stdin
+ libc.src.stdio.stdout
+ libc.src.stdio.ungetc
+
+ # inttypes.h entrypoints
+ libc.src.inttypes.imaxabs
+ libc.src.inttypes.imaxdiv
+ libc.src.inttypes.strtoimax
+ libc.src.inttypes.strtoumax
+
+ # time.h entrypoints
+ libc.src.time.clock
+ libc.src.time.clock_gettime
+ libc.src.time.timespec_get
+ libc.src.time.nanosleep
+ libc.src.time.strftime
+ libc.src.time.strftime_l
+
+ # wchar.h entrypoints
+ libc.src.wchar.wcslen
+ libc.src.wchar.wctob
+
+ # locale.h entrypoints
+ libc.src.locale.localeconv
+ libc.src.locale.duplocale
+ libc.src.locale.freelocale
+ libc.src.locale.localeconv
+ libc.src.locale.newlocale
+ libc.src.locale.setlocale
+ libc.src.locale.uselocale
+)
+
+set(TARGET_LIBM_ENTRYPOINTS
+ # math.h entrypoints
+ libc.src.math.acos
+ libc.src.math.acosf
+ libc.src.math.acosh
+ libc.src.math.acoshf
+ libc.src.math.asin
+ libc.src.math.asinf
+ libc.src.math.asinh
+ libc.src.math.asinhf
+ libc.src.math.atan
+ libc.src.math.atan2
+ libc.src.math.atan2f
+ libc.src.math.atan2l
+ libc.src.math.atanf
+ libc.src.math.atanh
+ libc.src.math.atanhf
+ libc.src.math.canonicalize
+ libc.src.math.canonicalizef
+ libc.src.math.canonicalizel
+ libc.src.math.cbrt
+ libc.src.math.cbrtf
+ libc.src.math.ceil
+ libc.src.math.ceilf
+ libc.src.math.ceill
+ libc.src.math.copysign
+ libc.src.math.copysignf
+ libc.src.math.copysignl
+ libc.src.math.cos
+ libc.src.math.cosf
+ libc.src.math.cosh
+ libc.src.math.coshf
+ libc.src.math.cospif
+ libc.src.math.ddivl
+ libc.src.math.dfmal
+ libc.src.math.dmull
+ libc.src.math.dsqrtl
+ libc.src.math.erf
+ libc.src.math.erff
+ libc.src.math.exp
+ libc.src.math.exp10
+ libc.src.math.exp10f
+ libc.src.math.exp2
+ libc.src.math.exp2f
+ libc.src.math.exp2m1f
+ libc.src.math.expf
+ libc.src.math.expm1
+ libc.src.math.expm1f
+ libc.src.math.fabs
+ libc.src.math.fabsf
+ libc.src.math.fabsl
+ libc.src.math.fadd
+ libc.src.math.faddl
+ libc.src.math.fdim
+ libc.src.math.fdimf
+ libc.src.math.fdiml
+ libc.src.math.fdiv
+ libc.src.math.fdivl
+ libc.src.math.ffma
+ libc.src.math.ffmal
+ libc.src.math.floor
+ libc.src.math.floorf
+ libc.src.math.floorl
+ libc.src.math.fma
+ libc.src.math.fmaf
+ libc.src.math.fmax
+ libc.src.math.fmaxf
+ libc.src.math.fmaximum
+ libc.src.math.fmaximumf
+ libc.src.math.fmaximuml
+ libc.src.math.fmaximum_mag
+ libc.src.math.fmaximum_magf
+ libc.src.math.fmaximum_magl
+ libc.src.math.fmaximum_mag_num
+ libc.src.math.fmaximum_mag_numf
+ libc.src.math.fmaximum_mag_numl
+ libc.src.math.fmaximum_num
+ libc.src.math.fmaximum_numf
+ libc.src.math.fmaximum_numl
+ libc.src.math.fmaxl
+ libc.src.math.fmin
+ libc.src.math.fminf
+ libc.src.math.fminimum
+ libc.src.math.fminimumf
+ libc.src.math.fminimuml
+ libc.src.math.fminimum_mag
+ libc.src.math.fminimum_magf
+ libc.src.math.fminimum_magl
+ libc.src.math.fminimum_mag_num
+ libc.src.math.fminimum_mag_numf
+ libc.src.math.fminimum_mag_numl
+ libc.src.math.fminimum_num
+ libc.src.math.fminimum_numf
+ libc.src.math.fminimum_numl
+ libc.src.math.fminl
+ libc.src.math.fmod
+ libc.src.math.fmodf
+ libc.src.math.fmodl
+ libc.src.math.fmul
+ libc.src.math.fmull
+ libc.src.math.frexp
+ libc.src.math.frexpf
+ libc.src.math.frexpl
+ libc.src.math.fromfp
+ libc.src.math.fromfpf
+ libc.src.math.fromfpl
+ libc.src.math.fromfpx
+ libc.src.math.fromfpxf
+ libc.src.math.fromfpxl
+ libc.src.math.fsqrt
+ libc.src.math.fsqrtl
+ libc.src.math.fsub
+ libc.src.math.fsubl
+ libc.src.math.getpayload
+ libc.src.math.getpayloadf
+ libc.src.math.getpayloadl
+ libc.src.math.hypot
+ libc.src.math.hypotf
+ libc.src.math.ilogb
+ libc.src.math.ilogbf
+ libc.src.math.ilogbl
+ libc.src.math.isnan
+ libc.src.math.isnanf
+ libc.src.math.isnanl
+ libc.src.math.ldexp
+ libc.src.math.ldexpf
+ libc.src.math.ldexpl
+ libc.src.math.lgamma
+ libc.src.math.lgamma_r
+ libc.src.math.llogb
+ libc.src.math.llogbf
+ libc.src.math.llogbl
+ libc.src.math.llrint
+ libc.src.math.llrintf
+ libc.src.math.llrintl
+ libc.src.math.llround
+ libc.src.math.llroundf
+ libc.src.math.llroundl
+ libc.src.math.log
+ libc.src.math.log10
+ libc.src.math.log10f
+ libc.src.math.log1p
+ libc.src.math.log1pf
+ libc.src.math.log2
+ libc.src.math.log2f
+ libc.src.math.logb
+ libc.src.math.logbf
+ libc.src.math.logbl
+ libc.src.math.logf
+ libc.src.math.lrint
+ libc.src.math.lrintf
+ libc.src.math.lrintl
+ libc.src.math.lround
+ libc.src.math.lroundf
+ libc.src.math.lroundl
+ libc.src.math.modf
+ libc.src.math.modff
+ libc.src.math.modfl
+ libc.src.math.nan
+ libc.src.math.nanf
+ libc.src.math.nanl
+ libc.src.math.nearbyint
+ libc.src.math.nearbyintf
+ libc.src.math.nearbyintl
+ libc.src.math.nextafter
+ libc.src.math.nextafterf
+ libc.src.math.nextafterl
+ libc.src.math.nextdown
+ libc.src.math.nextdownf
+ libc.src.math.nextdownl
+ libc.src.math.nexttoward
+ libc.src.math.nexttowardf
+ libc.src.math.nexttowardl
+ libc.src.math.nextup
+ libc.src.math.nextupf
+ libc.src.math.nextupl
+ libc.src.math.pow
+ libc.src.math.powf
+ libc.src.math.powi
+ libc.src.math.powif
+ libc.src.math.remainder
+ libc.src.math.remainderf
+ libc.src.math.remainderl
+ libc.src.math.remquo
+ libc.src.math.remquof
+ libc.src.math.remquol
+ libc.src.math.rint
+ libc.src.math.rintf
+ libc.src.math.rintl
+ libc.src.math.roundeven
+ libc.src.math.roundevenf
+ libc.src.math.roundevenl
+ libc.src.math.round
+ libc.src.math.roundf
+ libc.src.math.roundl
+ libc.src.math.scalbln
+ libc.src.math.scalblnf
+ libc.src.math.scalblnl
+ libc.src.math.scalbn
+ libc.src.math.scalbnf
+ libc.src.math.scalbnl
+ libc.src.math.setpayload
+ libc.src.math.setpayloadf
+ libc.src.math.setpayloadl
+ libc.src.math.setpayloadsig
+ libc.src.math.setpayloadsigf
+ libc.src.math.setpayloadsigl
+ libc.src.math.sin
+ libc.src.math.sincos
+ libc.src.math.sincosf
+ libc.src.math.sinf
+ libc.src.math.sinh
+ libc.src.math.sinhf
+ libc.src.math.sinpif
+ libc.src.math.sqrt
+ libc.src.math.sqrtf
+ libc.src.math.sqrtl
+ libc.src.math.tan
+ libc.src.math.tanf
+ libc.src.math.tanh
+ libc.src.math.tanhf
+ libc.src.math.tgamma
+ libc.src.math.tgammaf
+ libc.src.math.totalorder
+ libc.src.math.totalorderf
+ libc.src.math.totalorderl
+ libc.src.math.totalordermag
+ libc.src.math.totalordermagf
+ libc.src.math.totalordermagl
+ libc.src.math.trunc
+ libc.src.math.truncf
+ libc.src.math.truncl
+ libc.src.math.ufromfp
+ libc.src.math.ufromfpf
+ libc.src.math.ufromfpl
+ libc.src.math.ufromfpx
+ libc.src.math.ufromfpxf
+ libc.src.math.ufromfpxl
+)
+
+if(LIBC_TYPES_HAS_FLOAT16)
+ list(APPEND TARGET_LIBM_ENTRYPOINTS
+ # math.h C23 _Float16 entrypoints
+ libc.src.math.canonicalizef16
+ libc.src.math.ceilf16
+ libc.src.math.copysignf16
+ libc.src.math.coshf16
+ libc.src.math.exp10f16
+ libc.src.math.exp10m1f16
+ libc.src.math.exp2f16
+ libc.src.math.expf16
+ libc.src.math.f16add
+ libc.src.math.f16addf
+ libc.src.math.f16addl
+ libc.src.math.f16div
+ libc.src.math.f16divf
+ libc.src.math.f16divl
+ libc.src.math.f16fma
+ libc.src.math.f16fmaf
+ libc.src.math.f16fmal
+ libc.src.math.f16mul
+ libc.src.math.f16mulf
+ libc.src.math.f16mull
+ libc.src.math.f16sqrt
+ libc.src.math.f16sqrtf
+ libc.src.math.f16sqrtl
+ libc.src.math.f16sub
+ libc.src.math.f16subf
+ libc.src.math.f16subl
+ libc.src.math.fabsf16
+ libc.src.math.fdimf16
+ libc.src.math.floorf16
+ libc.src.math.fmaxf16
+ libc.src.math.fmaximum_mag_numf16
+ libc.src.math.fmaximum_magf16
+ libc.src.math.fmaximum_numf16
+ libc.src.math.fmaximumf16
+ libc.src.math.fminf16
+ libc.src.math.fminimum_mag_numf16
+ libc.src.math.fminimum_magf16
+ libc.src.math.fminimum_numf16
+ libc.src.math.fminimumf16
+ libc.src.math.fmodf16
+ libc.src.math.frexpf16
+ libc.src.math.fromfpf16
+ libc.src.math.fromfpxf16
+ libc.src.math.getpayloadf16
+ libc.src.math.ilogbf16
+ libc.src.math.ldexpf16
+ libc.src.math.llogbf16
+ libc.src.math.llrintf16
+ libc.src.math.llroundf16
+ libc.src.math.log10f16
+ libc.src.math.log2f16
+ libc.src.math.logbf16
+ libc.src.math.logf16
+ libc.src.math.lrintf16
+ libc.src.math.lroundf16
+ libc.src.math.modff16
+ libc.src.math.nanf16
+ libc.src.math.nearbyintf16
+ libc.src.math.nextafterf16
+ libc.src.math.nextdownf16
+ libc.src.math.nexttowardf16
+ libc.src.math.nextupf16
+ libc.src.math.remainderf16
+ libc.src.math.remquof16
+ libc.src.math.rintf16
+ libc.src.math.roundevenf16
+ libc.src.math.roundf16
+ libc.src.math.scalblnf16
+ libc.src.math.scalbnf16
+ libc.src.math.setpayloadf16
+ libc.src.math.setpayloadsigf16
+ libc.src.math.sinhf16
+ libc.src.math.sqrtf16
+ libc.src.math.tanhf16
+ libc.src.math.totalorderf16
+ libc.src.math.totalordermagf16
+ libc.src.math.truncf16
+ libc.src.math.ufromfpf16
+ libc.src.math.ufromfpxf16
+ )
+endif()
+
+set(TARGET_LLVMLIBC_ENTRYPOINTS
+ ${TARGET_LIBC_ENTRYPOINTS}
+ ${TARGET_LIBM_ENTRYPOINTS}
+)
diff --git a/libc/config/gpu/spirv64/headers.txt b/libc/config/gpu/spirv64/headers.txt
new file mode 100644
index 0000000000000..fa8ad7c11ba8b
--- /dev/null
+++ b/libc/config/gpu/spirv64/headers.txt
@@ -0,0 +1,21 @@
+set(TARGET_PUBLIC_HEADERS
+ libc.include.assert
+ libc.include.ctype
+ libc.include.string
+ libc.include.strings
+ libc.include.signal
+ libc.include.float
+ libc.include.stdint
+ libc.include.inttypes
+ libc.include.limits
+ libc.include.math
+ libc.include.fenv
+ libc.include.time
+ libc.include.errno
+ libc.include.stdlib
+ libc.include.stdio
+ libc.include.wchar
+ libc.include.uchar
+ libc.include.features
+ libc.include.locale
+)
diff --git a/libc/src/math/spirv64/CMakeLists.txt b/libc/src/math/spirv64/CMakeLists.txt
new file mode 100644
index 0000000000000..23fd3169fddfe
--- /dev/null
+++ b/libc/src/math/spirv64/CMakeLists.txt
@@ -0,0 +1,541 @@
+# 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)
+ 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()
+
+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(
+ 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
+)
+
+add_entrypoint_object(
+ frexp
+ SRCS
+ frexp.cpp
+ HDRS
+ ../frexp.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ frexpf
+ SRCS
+ frexpf.cpp
+ HDRS
+ ../frexpf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ scalbn
+ SRCS
+ scalbn.cpp
+ HDRS
+ ../scalbn.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ scalbnf
+ SRCS
+ scalbnf.cpp
+ HDRS
+ ../scalbnf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ ldexp
+ SRCS
+ ldexp.cpp
+ HDRS
+ ../ldexp.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+add_entrypoint_object(
+ ldexpf
+ SRCS
+ ldexpf.cpp
+ HDRS
+ ../ldexpf.h
+ COMPILE_OPTIONS
+ -O2
+)
+
+# 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
+ VENDOR
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+)
+
+add_entrypoint_object(
+ acosh
+ SRCS
+ acosh.cpp
+ HDRS
+ ../acosh.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(
+ 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(
+ atanh
+ SRCS
+ atanh.cpp
+ HDRS
+ ../atanh.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(
+ erf
+ SRCS
+ erf.cpp
+ HDRS
+ ../erf.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ powi
+ SRCS
+ powi.cpp
+ HDRS
+ ../powi.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ powif
+ SRCS
+ powif.cpp
+ HDRS
+ ../powif.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(
+ tanh
+ SRCS
+ tanh.cpp
+ HDRS
+ ../tanh.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(
+ lgamma
+ SRCS
+ lgamma.cpp
+ HDRS
+ ../lgamma.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ lgamma_r
+ SRCS
+ lgamma_r.cpp
+ HDRS
+ ../lgamma_r.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
diff --git a/libc/src/math/spirv64/acos.cpp b/libc/src/math/spirv64/acos.cpp
new file mode 100644
index 0000000000000..de870f207326e
--- /dev/null
+++ b/libc/src/math/spirv64/acos.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, acos, (double x)) { return __ocml_acos_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/acosf.cpp b/libc/src/math/spirv64/acosf.cpp
new file mode 100644
index 0000000000000..0a72a70a3ee9d
--- /dev/null
+++ b/libc/src/math/spirv64/acosf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, acosf, (float x)) { return __ocml_acos_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/acosh.cpp b/libc/src/math/spirv64/acosh.cpp
new file mode 100644
index 0000000000000..15c9734499832
--- /dev/null
+++ b/libc/src/math/spirv64/acosh.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, acosh, (double x)) { return __ocml_acosh_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/acoshf.cpp b/libc/src/math/spirv64/acoshf.cpp
new file mode 100644
index 0000000000000..79e71b02e7198
--- /dev/null
+++ b/libc/src/math/spirv64/acoshf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, acoshf, (float x)) { return __ocml_acosh_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/asin.cpp b/libc/src/math/spirv64/asin.cpp
new file mode 100644
index 0000000000000..a79641e1977cb
--- /dev/null
+++ b/libc/src/math/spirv64/asin.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, asin, (double x)) { return __ocml_asin_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/asinf.cpp b/libc/src/math/spirv64/asinf.cpp
new file mode 100644
index 0000000000000..e70944a4d9890
--- /dev/null
+++ b/libc/src/math/spirv64/asinf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, asinf, (float x)) { return __ocml_asin_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/asinh.cpp b/libc/src/math/spirv64/asinh.cpp
new file mode 100644
index 0000000000000..642368592a99b
--- /dev/null
+++ b/libc/src/math/spirv64/asinh.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, asinh, (double x)) { return __ocml_asinh_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/asinhf.cpp b/libc/src/math/spirv64/asinhf.cpp
new file mode 100644
index 0000000000000..bafa77f946000
--- /dev/null
+++ b/libc/src/math/spirv64/asinhf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, asinhf, (float x)) { return __ocml_asinh_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/atan.cpp b/libc/src/math/spirv64/atan.cpp
new file mode 100644
index 0000000000000..49941e97096f7
--- /dev/null
+++ b/libc/src/math/spirv64/atan.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, atan, (double x)) { return __ocml_atan_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/atan2.cpp b/libc/src/math/spirv64/atan2.cpp
new file mode 100644
index 0000000000000..f5907504cb364
--- /dev/null
+++ b/libc/src/math/spirv64/atan2.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, atan2, (double x, double y)) {
+ return __ocml_atan2_f64(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/atan2f.cpp b/libc/src/math/spirv64/atan2f.cpp
new file mode 100644
index 0000000000000..736c77d1cbce9
--- /dev/null
+++ b/libc/src/math/spirv64/atan2f.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, atan2f, (float x, float y)) {
+ return __ocml_atan2_f32(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/atanf.cpp b/libc/src/math/spirv64/atanf.cpp
new file mode 100644
index 0000000000000..ab1837dd026b7
--- /dev/null
+++ b/libc/src/math/spirv64/atanf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, atanf, (float x)) { return __ocml_atan_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/atanh.cpp b/libc/src/math/spirv64/atanh.cpp
new file mode 100644
index 0000000000000..091c155190f94
--- /dev/null
+++ b/libc/src/math/spirv64/atanh.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, atanh, (double x)) { return __ocml_atanh_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/atanhf.cpp b/libc/src/math/spirv64/atanhf.cpp
new file mode 100644
index 0000000000000..fa9cf39383963
--- /dev/null
+++ b/libc/src/math/spirv64/atanhf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, atanhf, (float x)) { return __ocml_atanh_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/ceil.cpp b/libc/src/math/spirv64/ceil.cpp
new file mode 100644
index 0000000000000..8834c7b560a1b
--- /dev/null
+++ b/libc/src/math/spirv64/ceil.cpp
@@ -0,0 +1,17 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, ceil, (double x)) { return __builtin_ceil(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/ceilf.cpp b/libc/src/math/spirv64/ceilf.cpp
new file mode 100644
index 0000000000000..5d26a30c849cd
--- /dev/null
+++ b/libc/src/math/spirv64/ceilf.cpp
@@ -0,0 +1,17 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, ceilf, (float x)) { return __builtin_ceilf(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/copysign.cpp b/libc/src/math/spirv64/copysign.cpp
new file mode 100644
index 0000000000000..06ef36fb3595f
--- /dev/null
+++ b/libc/src/math/spirv64/copysign.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, copysign, (double x, double y)) {
+ return __builtin_copysign(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/copysignf.cpp b/libc/src/math/spirv64/copysignf.cpp
new file mode 100644
index 0000000000000..aea94f3577d8f
--- /dev/null
+++ b/libc/src/math/spirv64/copysignf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, copysignf, (float x, float y)) {
+ return __builtin_copysignf(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/cos.cpp b/libc/src/math/spirv64/cos.cpp
new file mode 100644
index 0000000000000..a4d4c94dd7ac7
--- /dev/null
+++ b/libc/src/math/spirv64/cos.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, cos, (double x)) { return __ocml_cos_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/cosf.cpp b/libc/src/math/spirv64/cosf.cpp
new file mode 100644
index 0000000000000..99ec1185a8346
--- /dev/null
+++ b/libc/src/math/spirv64/cosf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, cosf, (float x)) { return __ocml_cos_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/cosh.cpp b/libc/src/math/spirv64/cosh.cpp
new file mode 100644
index 0000000000000..d94d7af251261
--- /dev/null
+++ b/libc/src/math/spirv64/cosh.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, cosh, (double x)) { return __ocml_cosh_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/coshf.cpp b/libc/src/math/spirv64/coshf.cpp
new file mode 100644
index 0000000000000..5b641be27818a
--- /dev/null
+++ b/libc/src/math/spirv64/coshf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, coshf, (float x)) { return __ocml_cosh_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/declarations.h b/libc/src/math/spirv64/declarations.h
new file mode 100644
index 0000000000000..88e2201521b67
--- /dev/null
+++ b/libc/src/math/spirv64/declarations.h
@@ -0,0 +1,91 @@
+//===-- AMDGPU specific declarations for math support ---------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_LIBC_SRC_MATH_AMDGPU_DECLARATIONS_H
+#define LLVM_LIBC_SRC_MATH_AMDGPU_DECLARATIONS_H
+
+#include "platform.h"
+
+#include "src/__support/GPU/utils.h"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+extern "C" {
+float __ocml_acos_f32(float);
+double __ocml_acos_f64(double);
+float __ocml_acosh_f32(float);
+double __ocml_acosh_f64(double);
+float __ocml_asin_f32(float);
+double __ocml_asin_f64(double);
+float __ocml_asinh_f32(float);
+double __ocml_asinh_f64(double);
+float __ocml_atan_f32(float);
+double __ocml_atan_f64(double);
+float __ocml_atan2_f32(float, float);
+double __ocml_atan2_f64(double, double);
+float __ocml_atanh_f32(float);
+double __ocml_atanh_f64(double);
+float __ocml_cos_f32(float);
+double __ocml_cos_f64(double);
+float __ocml_cosh_f32(float);
+double __ocml_cosh_f64(double);
+float __ocml_erf_f32(float);
+double __ocml_erf_f64(double);
+float __ocml_exp_f32(float);
+double __ocml_exp_f64(double);
+float __ocml_exp2_f32(float);
+double __ocml_exp2_f64(double);
+float __ocml_exp10_f32(float);
+double __ocml_exp10_f64(double);
+double __ocml_exp2_f64(double);
+float __ocml_expm1_f32(float);
+double __ocml_expm1_f64(double);
+float __ocml_fdim_f32(float, float);
+double __ocml_fdim_f64(double, double);
+float __ocml_hypot_f32(float, float);
+double __ocml_hypot_f64(double, double);
+int __ocml_ilogb_f64(double);
+int __ocml_ilogb_f32(float);
+float __ocml_ldexp_f32(float, int);
+double __ocml_ldexp_f64(double, int);
+float __ocml_log10_f32(float);
+double __ocml_log10_f64(double);
+float __ocml_log1p_f32(float);
+double __ocml_log1p_f64(double);
+float __ocml_log2_f32(float);
+double __ocml_log2_f64(double);
+float __ocml_log_f32(float);
+double __ocml_log_f64(double);
+float __ocml_nextafter_f32(float, float);
+double __ocml_nextafter_f64(double, double);
+float __ocml_pow_f32(float, float);
+double __ocml_pow_f64(double, double);
+float __ocml_pown_f32(float, int);
+double __ocml_pown_f64(double, int);
+float __ocml_sin_f32(float);
+double __ocml_sin_f64(double);
+float __ocml_sincos_f32(float, float *);
+double __ocml_sincos_f64(double, double *);
+float __ocml_sinh_f32(float);
+double __ocml_sinh_f64(double);
+float __ocml_tan_f32(float);
+double __ocml_tan_f64(double);
+float __ocml_tanh_f32(float);
+double __ocml_tanh_f64(double);
+float __ocml_remquo_f32(float, float, gpu::Private<int> *);
+double __ocml_remquo_f64(double, double, gpu::Private<int> *);
+double __ocml_tgamma_f64(double);
+float __ocml_tgamma_f32(float);
+double __ocml_lgamma_f64(double);
+double __ocml_lgamma_r_f64(double, gpu::Private<int> *);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
+
+#endif // LLVM_LIBC_SRC_MATH_AMDGPU_DECLARATIONS_H
diff --git a/libc/src/math/spirv64/erf.cpp b/libc/src/math/spirv64/erf.cpp
new file mode 100644
index 0000000000000..07ae268ff2cdc
--- /dev/null
+++ b/libc/src/math/spirv64/erf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, erf, (double x)) { return __ocml_erf_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/erff.cpp b/libc/src/math/spirv64/erff.cpp
new file mode 100644
index 0000000000000..a4b7b275cfcc2
--- /dev/null
+++ b/libc/src/math/spirv64/erff.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, erff, (float x)) { return __ocml_erf_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/exp.cpp b/libc/src/math/spirv64/exp.cpp
new file mode 100644
index 0000000000000..dae79bec5599b
--- /dev/null
+++ b/libc/src/math/spirv64/exp.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, exp, (double x)) { return __ocml_exp_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/exp10.cpp b/libc/src/math/spirv64/exp10.cpp
new file mode 100644
index 0000000000000..f13d218ca88c4
--- /dev/null
+++ b/libc/src/math/spirv64/exp10.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, exp10, (double x)) { return __ocml_exp10_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/exp10f.cpp b/libc/src/math/spirv64/exp10f.cpp
new file mode 100644
index 0000000000000..883e734e06877
--- /dev/null
+++ b/libc/src/math/spirv64/exp10f.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, exp10f, (float x)) { return __ocml_exp10_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/exp2.cpp b/libc/src/math/spirv64/exp2.cpp
new file mode 100644
index 0000000000000..fb336cf48f157
--- /dev/null
+++ b/libc/src/math/spirv64/exp2.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, exp2, (double x)) { return __ocml_exp2_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/exp2f.cpp b/libc/src/math/spirv64/exp2f.cpp
new file mode 100644
index 0000000000000..77b4a9c3b86e7
--- /dev/null
+++ b/libc/src/math/spirv64/exp2f.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, exp2f, (float x)) { return __ocml_exp2_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/expf.cpp b/libc/src/math/spirv64/expf.cpp
new file mode 100644
index 0000000000000..6c44aad164bfd
--- /dev/null
+++ b/libc/src/math/spirv64/expf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, expf, (float x)) { return __ocml_exp_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/expm1.cpp b/libc/src/math/spirv64/expm1.cpp
new file mode 100644
index 0000000000000..df3643f9c78b5
--- /dev/null
+++ b/libc/src/math/spirv64/expm1.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, expm1, (double x)) { return __ocml_expm1_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/expm1f.cpp b/libc/src/math/spirv64/expm1f.cpp
new file mode 100644
index 0000000000000..2409997f479c3
--- /dev/null
+++ b/libc/src/math/spirv64/expm1f.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, expm1f, (float x)) { return __ocml_expm1_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/fabs.cpp b/libc/src/math/spirv64/fabs.cpp
new file mode 100644
index 0000000000000..bb37596b9d563
--- /dev/null
+++ b/libc/src/math/spirv64/fabs.cpp
@@ -0,0 +1,17 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, fabs, (double x)) { return __builtin_fabs(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/fabsf.cpp b/libc/src/math/spirv64/fabsf.cpp
new file mode 100644
index 0000000000000..2698618f3f1e1
--- /dev/null
+++ b/libc/src/math/spirv64/fabsf.cpp
@@ -0,0 +1,17 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, fabsf, (float x)) { return __builtin_fabsf(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/fdim.cpp b/libc/src/math/spirv64/fdim.cpp
new file mode 100644
index 0000000000000..8ade0b21503b1
--- /dev/null
+++ b/libc/src/math/spirv64/fdim.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, fdim, (double x, double y)) {
+ return __ocml_fdim_f64(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/fdimf.cpp b/libc/src/math/spirv64/fdimf.cpp
new file mode 100644
index 0000000000000..ed3855e5ce2d6
--- /dev/null
+++ b/libc/src/math/spirv64/fdimf.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, fdimf, (float x, float y)) {
+ return __ocml_fdim_f32(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/floor.cpp b/libc/src/math/spirv64/floor.cpp
new file mode 100644
index 0000000000000..564efa9a7da38
--- /dev/null
+++ b/libc/src/math/spirv64/floor.cpp
@@ -0,0 +1,17 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, floor, (double x)) { return __builtin_floor(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/floorf.cpp b/libc/src/math/spirv64/floorf.cpp
new file mode 100644
index 0000000000000..6717c8f60c992
--- /dev/null
+++ b/libc/src/math/spirv64/floorf.cpp
@@ -0,0 +1,17 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, floorf, (float x)) { return __builtin_floorf(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/fma.cpp b/libc/src/math/spirv64/fma.cpp
new file mode 100644
index 0000000000000..c4a117e42a3aa
--- /dev/null
+++ b/libc/src/math/spirv64/fma.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, fma, (double x, double y, double z)) {
+ return __builtin_fma(x, y, z);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/fmaf.cpp b/libc/src/math/spirv64/fmaf.cpp
new file mode 100644
index 0000000000000..c088bd5b30fea
--- /dev/null
+++ b/libc/src/math/spirv64/fmaf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, fmaf, (float x, float y, float z)) {
+ return __builtin_fmaf(x, y, z);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/fmax.cpp b/libc/src/math/spirv64/fmax.cpp
new file mode 100644
index 0000000000000..474019733598f
--- /dev/null
+++ b/libc/src/math/spirv64/fmax.cpp
@@ -0,0 +1,22 @@
+//===-- 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/CPP/bit.h"
+#include "src/__support/common.h"
+#include "src/__support/macros/config.h"
+#include "src/__support/macros/optimization.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, fmax, (double x, double y)) {
+ return __builtin_fmax(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/fmaxf.cpp b/libc/src/math/spirv64/fmaxf.cpp
new file mode 100644
index 0000000000000..59ee8259c025a
--- /dev/null
+++ b/libc/src/math/spirv64/fmaxf.cpp
@@ -0,0 +1,20 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, fmaxf, (float x, float y)) {
+ return __builtin_fmaxf(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/fmin.cpp b/libc/src/math/spirv64/fmin.cpp
new file mode 100644
index 0000000000000..694eb664b599b
--- /dev/null
+++ b/libc/src/math/spirv64/fmin.cpp
@@ -0,0 +1,20 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, fmin, (double x, double y)) {
+ return __builtin_fmin(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/fminf.cpp b/libc/src/math/spirv64/fminf.cpp
new file mode 100644
index 0000000000000..2060b71b0841f
--- /dev/null
+++ b/libc/src/math/spirv64/fminf.cpp
@@ -0,0 +1,20 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, fminf, (float x, float y)) {
+ return __builtin_fminf(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/fmod.cpp b/libc/src/math/spirv64/fmod.cpp
new file mode 100644
index 0000000000000..49d19c4decb96
--- /dev/null
+++ b/libc/src/math/spirv64/fmod.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, fmod, (double x, double y)) {
+ return __builtin_fmod(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/fmodf.cpp b/libc/src/math/spirv64/fmodf.cpp
new file mode 100644
index 0000000000000..8fbcb0cc2ad9e
--- /dev/null
+++ b/libc/src/math/spirv64/fmodf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, fmodf, (float x, float y)) {
+ return __builtin_fmodf(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/frexp.cpp b/libc/src/math/spirv64/frexp.cpp
new file mode 100644
index 0000000000000..00e5187c0c719
--- /dev/null
+++ b/libc/src/math/spirv64/frexp.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, frexp, (double x, int *p)) {
+ return __builtin_frexp(x, p);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/frexpf.cpp b/libc/src/math/spirv64/frexpf.cpp
new file mode 100644
index 0000000000000..2799e542efae2
--- /dev/null
+++ b/libc/src/math/spirv64/frexpf.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, frexpf, (float x, int *p)) {
+ return __builtin_frexpf(x, p);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/hypot.cpp b/libc/src/math/spirv64/hypot.cpp
new file mode 100644
index 0000000000000..dcf1152a75160
--- /dev/null
+++ b/libc/src/math/spirv64/hypot.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, hypot, (double x, double y)) {
+ return __ocml_hypot_f64(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/hypotf.cpp b/libc/src/math/spirv64/hypotf.cpp
new file mode 100644
index 0000000000000..68ec659d12e5b
--- /dev/null
+++ b/libc/src/math/spirv64/hypotf.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, hypotf, (float x, float y)) {
+ return __ocml_hypot_f32(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/ilogb.cpp b/libc/src/math/spirv64/ilogb.cpp
new file mode 100644
index 0000000000000..37f24dfd8c39c
--- /dev/null
+++ b/libc/src/math/spirv64/ilogb.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(int, ilogb, (double x)) { return __ocml_ilogb_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/ilogbf.cpp b/libc/src/math/spirv64/ilogbf.cpp
new file mode 100644
index 0000000000000..56e74e1575f1d
--- /dev/null
+++ b/libc/src/math/spirv64/ilogbf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(int, ilogbf, (float x)) { return __ocml_ilogb_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/ldexp.cpp b/libc/src/math/spirv64/ldexp.cpp
new file mode 100644
index 0000000000000..393eabb84e61b
--- /dev/null
+++ b/libc/src/math/spirv64/ldexp.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, ldexp, (double x, int y)) {
+ return __builtin_ldexp(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/ldexpf.cpp b/libc/src/math/spirv64/ldexpf.cpp
new file mode 100644
index 0000000000000..970603dd170eb
--- /dev/null
+++ b/libc/src/math/spirv64/ldexpf.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, ldexpf, (float x, int y)) {
+ return __builtin_ldexpf(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/lgamma.cpp b/libc/src/math/spirv64/lgamma.cpp
new file mode 100644
index 0000000000000..acff4c76734e1
--- /dev/null
+++ b/libc/src/math/spirv64/lgamma.cpp
@@ -0,0 +1,19 @@
+//===-- Implementation of the lgamma 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/lgamma.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, lgamma, (double x)) { return __ocml_lgamma_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/lgamma_r.cpp b/libc/src/math/spirv64/lgamma_r.cpp
new file mode 100644
index 0000000000000..0a79988d66e30
--- /dev/null
+++ b/libc/src/math/spirv64/lgamma_r.cpp
@@ -0,0 +1,24 @@
+//===-- Implementation of the lgamma_r 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/lgamma_r.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, lgamma_r, (double x, int *signp)) {
+ int tmp = *signp;
+ double r = __ocml_lgamma_r_f64(x, (gpu::Private<int> *)&tmp);
+ *signp = tmp;
+ return r;
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/llrint.cpp b/libc/src/math/spirv64/llrint.cpp
new file mode 100644
index 0000000000000..21129fe80a0c1
--- /dev/null
+++ b/libc/src/math/spirv64/llrint.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(long long, llrint, (double x)) {
+ return static_cast<long long>(__builtin_rint(x));
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/llrintf.cpp b/libc/src/math/spirv64/llrintf.cpp
new file mode 100644
index 0000000000000..a6f9f43876010
--- /dev/null
+++ b/libc/src/math/spirv64/llrintf.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(long long, llrintf, (float x)) {
+ return static_cast<long long>(__builtin_rintf(x));
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/log.cpp b/libc/src/math/spirv64/log.cpp
new file mode 100644
index 0000000000000..bd01adfa08a4d
--- /dev/null
+++ b/libc/src/math/spirv64/log.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, log, (double x)) { return __ocml_log_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/log10.cpp b/libc/src/math/spirv64/log10.cpp
new file mode 100644
index 0000000000000..75957c9aeb506
--- /dev/null
+++ b/libc/src/math/spirv64/log10.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, log10, (double x)) { return __ocml_log10_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/log10f.cpp b/libc/src/math/spirv64/log10f.cpp
new file mode 100644
index 0000000000000..9c12d6b11bcf0
--- /dev/null
+++ b/libc/src/math/spirv64/log10f.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, log10f, (float x)) { return __ocml_log10_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/log1p.cpp b/libc/src/math/spirv64/log1p.cpp
new file mode 100644
index 0000000000000..fc275196830f6
--- /dev/null
+++ b/libc/src/math/spirv64/log1p.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, log1p, (double x)) { return __ocml_log1p_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/log1pf.cpp b/libc/src/math/spirv64/log1pf.cpp
new file mode 100644
index 0000000000000..b2d26fb1166fc
--- /dev/null
+++ b/libc/src/math/spirv64/log1pf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, log1pf, (float x)) { return __ocml_log1p_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/log2.cpp b/libc/src/math/spirv64/log2.cpp
new file mode 100644
index 0000000000000..73f34b65a32f7
--- /dev/null
+++ b/libc/src/math/spirv64/log2.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, log2, (double x)) { return __ocml_log2_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/log2f.cpp b/libc/src/math/spirv64/log2f.cpp
new file mode 100644
index 0000000000000..3b62edacfa656
--- /dev/null
+++ b/libc/src/math/spirv64/log2f.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, log2f, (float x)) { return __ocml_log2_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/logb.cpp b/libc/src/math/spirv64/logb.cpp
new file mode 100644
index 0000000000000..4b68e281be874
--- /dev/null
+++ b/libc/src/math/spirv64/logb.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, logb, (double x)) { return __ocml_logb_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/logbf.cpp b/libc/src/math/spirv64/logbf.cpp
new file mode 100644
index 0000000000000..bc7c462c7f927
--- /dev/null
+++ b/libc/src/math/spirv64/logbf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, logbf, (float x)) { return __ocml_logb_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/logf.cpp b/libc/src/math/spirv64/logf.cpp
new file mode 100644
index 0000000000000..17925672ef32f
--- /dev/null
+++ b/libc/src/math/spirv64/logf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, logf, (float x)) { return __ocml_log_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/lrint.cpp b/libc/src/math/spirv64/lrint.cpp
new file mode 100644
index 0000000000000..715b552575789
--- /dev/null
+++ b/libc/src/math/spirv64/lrint.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(long, lrint, (double x)) {
+ return static_cast<long>(__builtin_rint(x));
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/lrintf.cpp b/libc/src/math/spirv64/lrintf.cpp
new file mode 100644
index 0000000000000..387063803f27f
--- /dev/null
+++ b/libc/src/math/spirv64/lrintf.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(long, lrintf, (float x)) {
+ return static_cast<long>(__builtin_rintf(x));
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/nearbyint.cpp b/libc/src/math/spirv64/nearbyint.cpp
new file mode 100644
index 0000000000000..7d78c7241d023
--- /dev/null
+++ b/libc/src/math/spirv64/nearbyint.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, nearbyint, (double x)) {
+ return __builtin_nearbyint(x);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/nearbyintf.cpp b/libc/src/math/spirv64/nearbyintf.cpp
new file mode 100644
index 0000000000000..4bd20dc58fb89
--- /dev/null
+++ b/libc/src/math/spirv64/nearbyintf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, nearbyintf, (float x)) {
+ return __builtin_nearbyintf(x);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/nextafter.cpp b/libc/src/math/spirv64/nextafter.cpp
new file mode 100644
index 0000000000000..226b8a5a12547
--- /dev/null
+++ b/libc/src/math/spirv64/nextafter.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, nextafter, (double x, double y)) {
+ return __ocml_nextafter_f64(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/nextafterf.cpp b/libc/src/math/spirv64/nextafterf.cpp
new file mode 100644
index 0000000000000..7bed2c17b6e12
--- /dev/null
+++ b/libc/src/math/spirv64/nextafterf.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, nextafterf, (float x, float y)) {
+ return __ocml_nextafter_f32(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/platform.h b/libc/src/math/spirv64/platform.h
new file mode 100644
index 0000000000000..472a9830e8bb3
--- /dev/null
+++ b/libc/src/math/spirv64/platform.h
@@ -0,0 +1,55 @@
+//===-- AMDGPU specific platform 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_AMDGPU_PLATFORM_H
+#define LLVM_LIBC_SRC_MATH_AMDGPU_PLATFORM_H
+
+#include "src/__support/macros/attributes.h"
+#include "src/__support/macros/config.h"
+
+#include <stdint.h>
+
+namespace LIBC_NAMESPACE_DECL {
+
+// The ROCm device library uses control globals to alter codegen for the
+// different targets. To avoid needing to link them in manually we simply
+// define them here.
+extern "C" {
+
+// Disable unsafe math optimizations in the implementation.
+extern const LIBC_INLINE_VAR uint8_t __oclc_unsafe_math_opt = 0;
+
+// Disable denormalization at zero optimizations in the implementation.
+extern const LIBC_INLINE_VAR uint8_t __oclc_daz_opt = 0;
+
+// Disable rounding optimizations for 32-bit square roots.
+extern const LIBC_INLINE_VAR uint8_t __oclc_correctly_rounded_sqrt32 = 1;
+
+// Disable finite math optimizations.
+extern const LIBC_INLINE_VAR uint8_t __oclc_finite_only_opt = 0;
+
+// Set the ISA value to a high enough value that the ROCm device library math
+// functions will assume we have fast FMA operations among other features. This
+// is determined to be safe on all targets by looking at the source code.
+// https://github.com/ROCm/ROCm-Device-Libs/blob/amd-stg-open/ocml/src/opts.h
+extern const LIBC_INLINE_VAR uint32_t __oclc_ISA_version = 9000;
+}
+
+// These aliases cause clang to emit the control constants with ODR linkage.
+// This allows us to link against the symbols without preventing them from being
+// optimized out or causing symbol collisions.
+[[gnu::alias("__oclc_unsafe_math_opt")]] const uint8_t __oclc_unsafe_math_opt__;
+[[gnu::alias("__oclc_daz_opt")]] const uint8_t __oclc_daz_opt__;
+[[gnu::alias("__oclc_correctly_rounded_sqrt32")]] const uint8_t
+ __oclc_correctly_rounded_sqrt32__;
+[[gnu::alias("__oclc_finite_only_opt")]] const uint8_t __oclc_finite_only_opt__;
+[[gnu::alias("__oclc_ISA_version")]] const uint32_t __oclc_ISA_version__;
+
+} // namespace LIBC_NAMESPACE_DECL
+
+#endif // LLVM_LIBC_SRC_MATH_AMDGPU_PLATFORM_H
diff --git a/libc/src/math/spirv64/powf.cpp b/libc/src/math/spirv64/powf.cpp
new file mode 100644
index 0000000000000..6931934d4c11e
--- /dev/null
+++ b/libc/src/math/spirv64/powf.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, powf, (float x, float y)) {
+ return __ocml_pow_f32(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/powi.cpp b/libc/src/math/spirv64/powi.cpp
new file mode 100644
index 0000000000000..6b31b471e4b20
--- /dev/null
+++ b/libc/src/math/spirv64/powi.cpp
@@ -0,0 +1,21 @@
+//===-- Implementation of the powi 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/powi.h"
+#include "src/__support/common.h"
+#include "src/__support/macros/config.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, powi, (double x, int y)) {
+ return __ocml_pown_f64(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/powif.cpp b/libc/src/math/spirv64/powif.cpp
new file mode 100644
index 0000000000000..94f8a91d9c03b
--- /dev/null
+++ b/libc/src/math/spirv64/powif.cpp
@@ -0,0 +1,21 @@
+//===-- Implementation of the powi 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/powif.h"
+#include "src/__support/common.h"
+#include "src/__support/macros/config.h"
+
+#include "declarations.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, powif, (float x, int y)) {
+ return __ocml_pown_f32(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/remainder.cpp b/libc/src/math/spirv64/remainder.cpp
new file mode 100644
index 0000000000000..9027204312e00
--- /dev/null
+++ b/libc/src/math/spirv64/remainder.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, remainder, (double x, double y)) {
+ return __builtin_remainder(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/remainderf.cpp b/libc/src/math/spirv64/remainderf.cpp
new file mode 100644
index 0000000000000..50df3b2ce25c1
--- /dev/null
+++ b/libc/src/math/spirv64/remainderf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, remainderf, (float x, float y)) {
+ return __builtin_remainderf(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/remquo.cpp b/libc/src/math/spirv64/remquo.cpp
new file mode 100644
index 0000000000000..42c908e91c6c0
--- /dev/null
+++ b/libc/src/math/spirv64/remquo.cpp
@@ -0,0 +1,24 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+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_DECL
diff --git a/libc/src/math/spirv64/remquof.cpp b/libc/src/math/spirv64/remquof.cpp
new file mode 100644
index 0000000000000..854d3bfe842fd
--- /dev/null
+++ b/libc/src/math/spirv64/remquof.cpp
@@ -0,0 +1,24 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+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_DECL
diff --git a/libc/src/math/spirv64/rint.cpp b/libc/src/math/spirv64/rint.cpp
new file mode 100644
index 0000000000000..ac6837a4abc37
--- /dev/null
+++ b/libc/src/math/spirv64/rint.cpp
@@ -0,0 +1,17 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, rint, (double x)) { return __builtin_rint(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/rintf.cpp b/libc/src/math/spirv64/rintf.cpp
new file mode 100644
index 0000000000000..94093471a8d92
--- /dev/null
+++ b/libc/src/math/spirv64/rintf.cpp
@@ -0,0 +1,17 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, rintf, (float x)) { return __builtin_rintf(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/round.cpp b/libc/src/math/spirv64/round.cpp
new file mode 100644
index 0000000000000..0d2765f2e959f
--- /dev/null
+++ b/libc/src/math/spirv64/round.cpp
@@ -0,0 +1,17 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, round, (double x)) { return __builtin_round(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/roundf.cpp b/libc/src/math/spirv64/roundf.cpp
new file mode 100644
index 0000000000000..86e8ba3ac83d6
--- /dev/null
+++ b/libc/src/math/spirv64/roundf.cpp
@@ -0,0 +1,17 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, roundf, (float x)) { return __builtin_roundf(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/scalbn.cpp b/libc/src/math/spirv64/scalbn.cpp
new file mode 100644
index 0000000000000..05bbbc69b47a1
--- /dev/null
+++ b/libc/src/math/spirv64/scalbn.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, scalbn, (double x, int y)) {
+ return __builtin_amdgcn_ldexp(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/scalbnf.cpp b/libc/src/math/spirv64/scalbnf.cpp
new file mode 100644
index 0000000000000..f0e9e475674c3
--- /dev/null
+++ b/libc/src/math/spirv64/scalbnf.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, scalbnf, (float x, int y)) {
+ return __builtin_amdgcn_ldexpf(x, y);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/sin.cpp b/libc/src/math/spirv64/sin.cpp
new file mode 100644
index 0000000000000..f3d88af009ce4
--- /dev/null
+++ b/libc/src/math/spirv64/sin.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, sin, (double x)) { return __ocml_sin_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/sincos.cpp b/libc/src/math/spirv64/sincos.cpp
new file mode 100644
index 0000000000000..304ac0c7c6332
--- /dev/null
+++ b/libc/src/math/spirv64/sincos.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(void, sincos, (double x, double *sinptr, double *cosptr)) {
+ *sinptr = __ocml_sincos_f64(x, cosptr);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/sincosf.cpp b/libc/src/math/spirv64/sincosf.cpp
new file mode 100644
index 0000000000000..1c4e9c6a2463f
--- /dev/null
+++ b/libc/src/math/spirv64/sincosf.cpp
@@ -0,0 +1,21 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(void, sincosf, (float x, float *sinptr, float *cosptr)) {
+ *sinptr = __ocml_sincos_f32(x, cosptr);
+}
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/sinf.cpp b/libc/src/math/spirv64/sinf.cpp
new file mode 100644
index 0000000000000..c6d64a63c755d
--- /dev/null
+++ b/libc/src/math/spirv64/sinf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, sinf, (float x)) { return __ocml_sin_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/sinh.cpp b/libc/src/math/spirv64/sinh.cpp
new file mode 100644
index 0000000000000..26314f46d0e97
--- /dev/null
+++ b/libc/src/math/spirv64/sinh.cpp
@@ -0,0 +1,19 @@
+//===-- 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.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/sinh.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, sinh, (double x)) { return __ocml_sinh_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/sinhf.cpp b/libc/src/math/spirv64/sinhf.cpp
new file mode 100644
index 0000000000000..a4eb8e1a65283
--- /dev/null
+++ b/libc/src/math/spirv64/sinhf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, sinhf, (float x)) { return __ocml_sinh_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/sqrt.cpp b/libc/src/math/spirv64/sqrt.cpp
new file mode 100644
index 0000000000000..ed83b6e5c6cae
--- /dev/null
+++ b/libc/src/math/spirv64/sqrt.cpp
@@ -0,0 +1,17 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, sqrt, (double x)) { return __builtin_sqrt(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/sqrtf.cpp b/libc/src/math/spirv64/sqrtf.cpp
new file mode 100644
index 0000000000000..851922c316452
--- /dev/null
+++ b/libc/src/math/spirv64/sqrtf.cpp
@@ -0,0 +1,17 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, sqrtf, (float x)) { return __builtin_sqrtf(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/tan.cpp b/libc/src/math/spirv64/tan.cpp
new file mode 100644
index 0000000000000..c946dc2a3db67
--- /dev/null
+++ b/libc/src/math/spirv64/tan.cpp
@@ -0,0 +1,19 @@
+//===-- 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.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/tan.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, tan, (double x)) { return __ocml_tan_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/tanf.cpp b/libc/src/math/spirv64/tanf.cpp
new file mode 100644
index 0000000000000..8c93fc4f62369
--- /dev/null
+++ b/libc/src/math/spirv64/tanf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, tanf, (float x)) { return __ocml_tan_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/tanh.cpp b/libc/src/math/spirv64/tanh.cpp
new file mode 100644
index 0000000000000..834353e122abe
--- /dev/null
+++ b/libc/src/math/spirv64/tanh.cpp
@@ -0,0 +1,19 @@
+//===-- 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.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/tanh.h"
+#include "src/__support/common.h"
+
+#include "declarations.h"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, tanh, (double x)) { return __ocml_tanh_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/tanhf.cpp b/libc/src/math/spirv64/tanhf.cpp
new file mode 100644
index 0000000000000..5029596a8e6fd
--- /dev/null
+++ b/libc/src/math/spirv64/tanhf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, tanhf, (float x)) { return __ocml_tanh_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/tgamma.cpp b/libc/src/math/spirv64/tgamma.cpp
new file mode 100644
index 0000000000000..485a6a3def432
--- /dev/null
+++ b/libc/src/math/spirv64/tgamma.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, tgamma, (double x)) { return __ocml_tgamma_f64(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/tgammaf.cpp b/libc/src/math/spirv64/tgammaf.cpp
new file mode 100644
index 0000000000000..e48a486c9c5e5
--- /dev/null
+++ b/libc/src/math/spirv64/tgammaf.cpp
@@ -0,0 +1,19 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, tgammaf, (float x)) { return __ocml_tgamma_f32(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/trunc.cpp b/libc/src/math/spirv64/trunc.cpp
new file mode 100644
index 0000000000000..f60caa2a71d78
--- /dev/null
+++ b/libc/src/math/spirv64/trunc.cpp
@@ -0,0 +1,17 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(double, trunc, (double x)) { return __builtin_trunc(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/src/math/spirv64/truncf.cpp b/libc/src/math/spirv64/truncf.cpp
new file mode 100644
index 0000000000000..a6c9b8f188f02
--- /dev/null
+++ b/libc/src/math/spirv64/truncf.cpp
@@ -0,0 +1,17 @@
+//===-- 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"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+LLVM_LIBC_FUNCTION(float, truncf, (float x)) { return __builtin_truncf(x); }
+
+} // namespace LIBC_NAMESPACE_DECL
diff --git a/libc/startup/gpu/spirv64/CMakeLists.txt b/libc/startup/gpu/spirv64/CMakeLists.txt
new file mode 100644
index 0000000000000..b67a5a2cc89fb
--- /dev/null
+++ b/libc/startup/gpu/spirv64/CMakeLists.txt
@@ -0,0 +1,15 @@
+add_startup_object(
+ crt1
+ SRC
+ start.cpp
+ DEPENDS
+ libc.config.app_h
+ libc.src.__support.RPC.rpc_client
+ libc.src.__support.GPU.utils
+ libc.src.stdlib.exit
+ libc.src.stdlib.atexit
+ COMPILE_OPTIONS
+ -ffreestanding # To avoid compiler warnings about calling the main function.
+ -fno-builtin
+)
+get_fq_target_name(crt1 fq_name)
diff --git a/libc/startup/gpu/spirv64/start.cpp b/libc/startup/gpu/spirv64/start.cpp
new file mode 100644
index 0000000000000..b1cb9f7be016b
--- /dev/null
+++ b/libc/startup/gpu/spirv64/start.cpp
@@ -0,0 +1,80 @@
+//===-- Implementation of crt for amdgpu ----------------------------------===//
+//
+// 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 "config/gpu/app.h"
+#include "src/__support/GPU/utils.h"
+#include "src/__support/RPC/rpc_client.h"
+#include "src/__support/macros/config.h"
+#include "src/stdlib/atexit.h"
+#include "src/stdlib/exit.h"
+
+extern "C" int main(int argc, char **argv, char **envp);
+
+namespace LIBC_NAMESPACE_DECL {
+
+// FIXME: Factor this out into common logic so we don't need to stub it here.
+void teardown_main_tls() {}
+
+DataEnvironment app;
+
+extern "C" uintptr_t __init_array_start[];
+extern "C" uintptr_t __init_array_end[];
+extern "C" uintptr_t __fini_array_start[];
+extern "C" uintptr_t __fini_array_end[];
+
+using InitCallback = void(int, char **, char **);
+using FiniCallback = void(void);
+
+static void call_init_array_callbacks(int argc, char **argv, char **env) {
+ size_t init_array_size = __init_array_end - __init_array_start;
+ for (size_t i = 0; i < init_array_size; ++i)
+ reinterpret_cast<InitCallback *>(__init_array_start[i])(argc, argv, env);
+}
+
+static void call_fini_array_callbacks() {
+ size_t fini_array_size = __fini_array_end - __fini_array_start;
+ for (size_t i = fini_array_size; i > 0; --i)
+ reinterpret_cast<FiniCallback *>(__fini_array_start[i - 1])();
+}
+
+} // namespace LIBC_NAMESPACE_DECL
+
+// spirv is rejecting amdgpu_kernel, kernel and __kernel
+// clang suggests it wants kernels marked with something from opencl
+
+extern "C" [[gnu::visibility("protected")/*, clang::amdgpu_kernel,
+ clang::amdgpu_flat_work_group_size(1, 1),
+ clang::amdgpu_max_num_work_groups(1)*/]]
+void
+_begin(int argc, char **argv, char **env) {
+ __atomic_store_n(&LIBC_NAMESPACE::app.env_ptr,
+ reinterpret_cast<uintptr_t *>(env), __ATOMIC_RELAXED);
+ // We want the fini array callbacks to be run after other atexit
+ // callbacks are run. So, we register them before running the init
+ // array callbacks as they can potentially register their own atexit
+ // callbacks.
+ LIBC_NAMESPACE::atexit(&LIBC_NAMESPACE::call_fini_array_callbacks);
+ LIBC_NAMESPACE::call_init_array_callbacks(argc, argv, env);
+}
+
+extern "C" [[gnu::visibility("protected")/*, clang::amdgpu_kernel*/]] void
+_start(int argc, char **argv, char **envp, int *ret) {
+ // Invoke the 'main' function with every active thread that the user launched
+ // the _start kernel with.
+ __atomic_fetch_or(ret, main(argc, argv, envp), __ATOMIC_RELAXED);
+}
+
+extern "C" [[gnu::visibility("protected")/*, clang::amdgpu_kernel,
+ clang::amdgpu_flat_work_group_size(1, 1),
+ clang::amdgpu_max_num_work_groups(1)*/]] void
+_end(int retval) {
+ // Only a single thread should call `exit` here, the rest should gracefully
+ // return from the kernel. This is so only one thread calls the destructors
+ // registred with 'atexit' above.
+ LIBC_NAMESPACE::exit(retval);
+}
diff --git a/lld/ELF/InputFiles.cpp b/lld/ELF/InputFiles.cpp
index d43de8ce6dfef..9549a1c3d48a6 100644
--- a/lld/ELF/InputFiles.cpp
+++ b/lld/ELF/InputFiles.cpp
@@ -1680,7 +1680,11 @@ static uint16_t getBitcodeMachineKind(Ctx &ctx, StringRef path,
case Triple::x86:
return t.isOSIAMCU() ? EM_IAMCU : EM_386;
case Triple::x86_64:
- return EM_X86_64;
+ return EM_X86_64;
+ case Triple::spirv64:
+ if (t.getVendor() == Triple::AMD)
+ return EM_AMDGPU;
+ LLVM_FALLTHROUGH;
default:
ErrAlways(ctx) << path
<< ": could not infer e_machine from bitcode target triple "
diff --git a/offload/cmake/caches/Offload.cmake b/offload/cmake/caches/Offload.cmake
index 6c74fc1d4b762..228cfae85784e 100644
--- a/offload/cmake/caches/Offload.cmake
+++ b/offload/cmake/caches/Offload.cmake
@@ -2,8 +2,10 @@ set(LLVM_ENABLE_PROJECTS "clang;clang-tools-extra;lld" CACHE STRING "")
set(LLVM_ENABLE_RUNTIMES "compiler-rt;openmp;offload" CACHE STRING "")
set(LLVM_ENABLE_PER_TARGET_RUNTIME_DIR ON CACHE BOOL "")
-set(LLVM_RUNTIME_TARGETS default;amdgcn-amd-amdhsa;nvptx64-nvidia-cuda CACHE STRING "")
+set(LLVM_RUNTIME_TARGETS default;amdgcn-amd-amdhsa;nvptx64-nvidia-cuda;spirv64-amd-amdhsa CACHE STRING "")
# set(RUNTIMES_nvptx64-nvidia-cuda_CACHE_FILES "${CMAKE_SOURCE_DIR}/../libcxx/cmake/caches/NVPTX.cmake" CACHE STRING "")
# set(RUNTIMES_amdgcn-amd-amdhsa_CACHE_FILES "${CMAKE_SOURCE_DIR}/../libcxx/cmake/caches/AMDGPU.cmake" CACHE STRING "")
set(RUNTIMES_nvptx64-nvidia-cuda_LLVM_ENABLE_RUNTIMES "compiler-rt;libc" CACHE STRING "")
set(RUNTIMES_amdgcn-amd-amdhsa_LLVM_ENABLE_RUNTIMES "compiler-rt;libc" CACHE STRING "")
+set(RUNTIMES_spirv64-amd-amdhsa_LLVM_ENABLE_RUNTIMES "libc" CACHE STRING "")
+set(RUNTIMES_spirv64-amd-amdhsa_LLVM_LIBC_FULL_BUILD ON CACHE BOOL "")
More information about the cfe-commits
mailing list