[clang] [libc] [SPIRV][libc] Introduce SPIR-V target for 'libc' (PR #137310)
Vyacheslav Levytskyy via cfe-commits
cfe-commits at lists.llvm.org
Mon Apr 28 09:37:42 PDT 2025
https://github.com/VyacheslavLevytskyy updated https://github.com/llvm/llvm-project/pull/137310
>From fa1b1b072051b077beb7342a5fcf10f378f44803 Mon Sep 17 00:00:00 2001
From: "Levytskyy, Vyacheslav" <vyacheslav.levytskyy at intel.com>
Date: Fri, 25 Apr 2025 03:51:37 -0700
Subject: [PATCH 1/2] introduce SPIR-V target for 'libc'
---
libc/benchmarks/gpu/timing/CMakeLists.txt | 2 +-
libc/benchmarks/gpu/timing/timing.h | 2 ++
libc/cmake/modules/LLVMLibCArchitectures.cmake | 3 ++-
libc/include/llvm-libc-macros/math-macros.h | 3 ++-
libc/include/llvm-libc-macros/signal-macros.h | 2 +-
libc/include/llvm-libc-macros/time-macros.h | 2 +-
libc/include/llvm-libc-types/fenv_t.h | 2 +-
libc/shared/rpc_util.h | 2 +-
libc/src/__support/macros/properties/architectures.h | 7 ++++++-
libc/src/__support/macros/properties/cpu_features.h | 5 +++--
10 files changed, 20 insertions(+), 10 deletions(-)
diff --git a/libc/benchmarks/gpu/timing/CMakeLists.txt b/libc/benchmarks/gpu/timing/CMakeLists.txt
index b6d84607aa607..94dd09d11e377 100644
--- a/libc/benchmarks/gpu/timing/CMakeLists.txt
+++ b/libc/benchmarks/gpu/timing/CMakeLists.txt
@@ -1,4 +1,4 @@
-foreach(target nvptx amdgpu)
+foreach(target nvptx amdgpu spirv64)
add_subdirectory(${target})
list(APPEND target_gpu_timing libc.benchmarks.gpu.timing.${target}.${target}_timing)
endforeach()
diff --git a/libc/benchmarks/gpu/timing/timing.h b/libc/benchmarks/gpu/timing/timing.h
index 2e098feb4b3a5..43337a46e7052 100644
--- a/libc/benchmarks/gpu/timing/timing.h
+++ b/libc/benchmarks/gpu/timing/timing.h
@@ -15,6 +15,8 @@
#include "amdgpu/timing.h"
#elif defined(LIBC_TARGET_ARCH_IS_NVPTX)
#include "nvptx/timing.h"
+#elif defined(LIBC_TARGET_ARCH_IS_SPIRV)
+#include "spirv64/timing.h"
#else
#error "unsupported platform"
#endif
diff --git a/libc/cmake/modules/LLVMLibCArchitectures.cmake b/libc/cmake/modules/LLVMLibCArchitectures.cmake
index 62f3a2e3bdb59..a5286ee9e96db 100644
--- a/libc/cmake/modules/LLVMLibCArchitectures.cmake
+++ b/libc/cmake/modules/LLVMLibCArchitectures.cmake
@@ -73,7 +73,8 @@ function(get_arch_and_system_from_triple triple arch_var sys_var)
# Setting OS name for GPU architectures.
list(GET triple_comps -1 gpu_target_sys)
- if(gpu_target_sys MATCHES "^amdhsa" OR gpu_target_sys MATCHES "^cuda")
+ if(gpu_target_sys MATCHES "^amdhsa" OR gpu_target_sys MATCHES "^cuda" OR
+ gpu_target_sys MATCHES "^opencl")
set(target_sys "gpu")
endif()
diff --git a/libc/include/llvm-libc-macros/math-macros.h b/libc/include/llvm-libc-macros/math-macros.h
index 2f05d7544666e..41d572492b223 100644
--- a/libc/include/llvm-libc-macros/math-macros.h
+++ b/libc/include/llvm-libc-macros/math-macros.h
@@ -42,7 +42,8 @@
#define FP_LLOGBNAN LONG_MAX
#endif
-#if defined(__NVPTX__) || defined(__AMDGPU__) || defined(__FAST_MATH__)
+#if defined(__NVPTX__) || defined(__AMDGPU__) || defined(__SPIRV__) || \
+ defined(__FAST_MATH__)
#define math_errhandling 0
#elif defined(__NO_MATH_ERRNO__)
#define math_errhandling (MATH_ERREXCEPT)
diff --git a/libc/include/llvm-libc-macros/signal-macros.h b/libc/include/llvm-libc-macros/signal-macros.h
index fbe929a0fea25..163c8742593b8 100644
--- a/libc/include/llvm-libc-macros/signal-macros.h
+++ b/libc/include/llvm-libc-macros/signal-macros.h
@@ -11,7 +11,7 @@
#if defined(__linux__)
#include "linux/signal-macros.h"
-#elif defined(__NVPTX__) || defined(__AMDGPU__)
+#elif defined(__NVPTX__) || defined(__AMDGPU__) || defined(__SPIRV__)
#include "gpu/signal-macros.h"
#endif
diff --git a/libc/include/llvm-libc-macros/time-macros.h b/libc/include/llvm-libc-macros/time-macros.h
index 445d8b3e837ed..8e1170ce8e408 100644
--- a/libc/include/llvm-libc-macros/time-macros.h
+++ b/libc/include/llvm-libc-macros/time-macros.h
@@ -1,7 +1,7 @@
#ifndef LLVM_LIBC_MACROS_TIME_MACROS_H
#define LLVM_LIBC_MACROS_TIME_MACROS_H
-#if defined(__AMDGPU__) || defined(__NVPTX__)
+#if defined(__AMDGPU__) || defined(__NVPTX__) || defined(__SPIRV__)
#include "gpu/time-macros.h"
#elif defined(__linux__)
#include "linux/time-macros.h"
diff --git a/libc/include/llvm-libc-types/fenv_t.h b/libc/include/llvm-libc-types/fenv_t.h
index c83f23894c0c8..2cfeff7b8a9f8 100644
--- a/libc/include/llvm-libc-types/fenv_t.h
+++ b/libc/include/llvm-libc-types/fenv_t.h
@@ -25,7 +25,7 @@ typedef struct {
} fenv_t;
#elif defined(__riscv)
typedef unsigned int fenv_t;
-#elif defined(__AMDGPU__) || defined(__NVPTX__)
+#elif defined(__AMDGPU__) || defined(__NVPTX__) || defined(__SPIRV__)
typedef struct {
unsigned int __fpc;
} fenv_t;
diff --git a/libc/shared/rpc_util.h b/libc/shared/rpc_util.h
index 687814b7ff2ae..1785920ec2504 100644
--- a/libc/shared/rpc_util.h
+++ b/libc/shared/rpc_util.h
@@ -12,7 +12,7 @@
#include <stddef.h>
#include <stdint.h>
-#if (defined(__NVPTX__) || defined(__AMDGPU__)) && \
+#if (defined(__NVPTX__) || defined(__AMDGPU__) || defined(__SPIRV__)) && \
!((defined(__CUDA__) && !defined(__CUDA_ARCH__)) || \
(defined(__HIP__) && !defined(__HIP_DEVICE_COMPILE__)))
#include <gpuintrin.h>
diff --git a/libc/src/__support/macros/properties/architectures.h b/libc/src/__support/macros/properties/architectures.h
index c88956ff41148..117c2eaca3ace 100644
--- a/libc/src/__support/macros/properties/architectures.h
+++ b/libc/src/__support/macros/properties/architectures.h
@@ -17,7 +17,12 @@
#define LIBC_TARGET_ARCH_IS_NVPTX
#endif
-#if defined(LIBC_TARGET_ARCH_IS_NVPTX) || defined(LIBC_TARGET_ARCH_IS_AMDGPU)
+#if defined(__SPIRV__)
+#define LIBC_TARGET_ARCH_IS_SPIRV
+#endif
+
+#if defined(LIBC_TARGET_ARCH_IS_NVPTX) || \
+ defined(LIBC_TARGET_ARCH_IS_AMDGPU) || defined(LIBC_TARGET_ARCH_IS_SPIRV)
#define LIBC_TARGET_ARCH_IS_GPU
#endif
diff --git a/libc/src/__support/macros/properties/cpu_features.h b/libc/src/__support/macros/properties/cpu_features.h
index 3677e1fc3275c..6f97b61596685 100644
--- a/libc/src/__support/macros/properties/cpu_features.h
+++ b/libc/src/__support/macros/properties/cpu_features.h
@@ -75,13 +75,14 @@
#endif // LIBC_TARGET_CPU_HAS_RISCV_FPU_DOUBLE
#endif // __riscv_flen
-#if defined(__NVPTX__) || defined(__AMDGPU__)
+#if defined(__NVPTX__) || defined(__AMDGPU__) || defined(__SPIRV__)
#define LIBC_TARGET_CPU_HAS_FPU_FLOAT
#define LIBC_TARGET_CPU_HAS_FPU_DOUBLE
#endif
#if defined(__ARM_FEATURE_FMA) || (defined(__AVX2__) && defined(__FMA__)) || \
- defined(__NVPTX__) || defined(__AMDGPU__) || defined(__LIBC_RISCV_USE_FMA)
+ defined(__NVPTX__) || defined(__AMDGPU__) || defined(__SPIRV__) || \
+ defined(__LIBC_RISCV_USE_FMA)
#define LIBC_TARGET_CPU_HAS_FMA
// Provide a more fine-grained control of FMA instruction for ARM targets.
#if defined(LIBC_TARGET_CPU_HAS_FPU_HALF)
>From c18d2e09c4cd841b6868e342e61e6f7ac34d27be Mon Sep 17 00:00:00 2001
From: "Levytskyy, Vyacheslav" <vyacheslav.levytskyy at intel.com>
Date: Mon, 28 Apr 2025 09:37:23 -0700
Subject: [PATCH 2/2] introduce SPIR-V target for 'libc'
---
clang/include/clang/Basic/BuiltinsSPIRV.td | 84 +++
clang/lib/Headers/CMakeLists.txt | 1 +
clang/lib/Headers/gpuintrin.h | 2 +
clang/lib/Headers/spirvintrin.h | 186 ++++++
.../gpu/timing/spirv64/CMakeLists.txt | 11 +
libc/benchmarks/gpu/timing/spirv64/timing.h | 152 +++++
.../cmake/modules/LLVMLibCArchitectures.cmake | 3 +-
libc/config/gpu/spirv64/config.json | 40 ++
libc/config/gpu/spirv64/entrypoints.txt | 610 ++++++++++++++++++
libc/config/gpu/spirv64/headers.txt | 21 +
libc/src/math/spirv64/CMakeLists.txt | 482 ++++++++++++++
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 | 21 +
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 | 20 +
libc/src/math/spirv64/fmaxf.cpp | 22 +
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 | 23 +
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 | 19 +
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 | 103 +++
libc/src/math/spirv64/powf.cpp | 19 +
libc/src/math/spirv64/powi.cpp | 19 +
libc/src/math/spirv64/powif.cpp | 19 +
libc/src/math/spirv64/remainder.cpp | 19 +
libc/src/math/spirv64/remainderf.cpp | 19 +
libc/src/math/spirv64/remquo.cpp | 21 +
libc/src/math/spirv64/remquof.cpp | 21 +
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 | 74 +++
118 files changed, 3785 insertions(+), 1 deletion(-)
create mode 100644 clang/lib/Headers/spirvintrin.h
create mode 100644 libc/benchmarks/gpu/timing/spirv64/CMakeLists.txt
create mode 100644 libc/benchmarks/gpu/timing/spirv64/timing.h
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/include/clang/Basic/BuiltinsSPIRV.td b/clang/include/clang/Basic/BuiltinsSPIRV.td
index 9f76d672cc7ce..4de95329a3824 100644
--- a/clang/include/clang/Basic/BuiltinsSPIRV.td
+++ b/clang/include/clang/Basic/BuiltinsSPIRV.td
@@ -31,3 +31,87 @@ def SPIRVSmoothStep : Builtin {
let Attributes = [NoThrow, Const, CustomTypeChecking];
let Prototype = "void(...)";
}
+
+def SPIRVGetNumWrkgrpX : Builtin {
+ let Spellings = ["__builtin_spirv_get_num_workgroups_x"];
+ let Attributes = [NoThrow, Const];
+ let Prototype = "void(...)";
+}
+
+def SPIRVGetNumWrkgrpY : Builtin {
+ let Spellings = ["__builtin_spirv_get_num_workgroups_y"];
+ let Attributes = [NoThrow, Const];
+ let Prototype = "void(...)";
+}
+
+def SPIRVGetNumWrkgrpZ : Builtin {
+ let Spellings = ["__builtin_spirv_get_num_workgroups_z"];
+ let Attributes = [NoThrow, Const];
+ let Prototype = "void(...)";
+}
+
+def SPIRVGetWrkgrpIdX : Builtin {
+ let Spellings = ["__builtin_spirv_get_workgroup_id_x"];
+ let Attributes = [NoThrow, Const];
+ let Prototype = "void(...)";
+}
+
+def SPIRVGetWrkgrpIdY : Builtin {
+ let Spellings = ["__builtin_spirv_get_workgroup_id_y"];
+ let Attributes = [NoThrow, Const];
+ let Prototype = "void(...)";
+}
+
+def SPIRVGetWrkgrpIdZ : Builtin {
+ let Spellings = ["__builtin_spirv_get_workgroup_id_z"];
+ let Attributes = [NoThrow, Const];
+ let Prototype = "void(...)";
+}
+
+def SPIRVGetWrkgrpSizeX : Builtin {
+ let Spellings = ["__builtin_spirv_workgroup_size_x"];
+ let Attributes = [NoThrow, Const];
+ let Prototype = "void(...)";
+}
+
+def SPIRVGetWrkgrpSizeY : Builtin {
+ let Spellings = ["__builtin_spirv_workgroup_size_y"];
+ let Attributes = [NoThrow, Const];
+ let Prototype = "void(...)";
+}
+
+def SPIRVGetWrkgrpSizeZ : Builtin {
+ let Spellings = ["__builtin_spirv_workgroup_size_z"];
+ let Attributes = [NoThrow, Const];
+ let Prototype = "void(...)";
+}
+
+def SPIRVGetWrkitemIdX : Builtin {
+ let Spellings = ["__builtin_spirv_workitem_id_x"];
+ let Attributes = [NoThrow, Const];
+ let Prototype = "void(...)";
+}
+
+def SPIRVGetWrkitemIdY : Builtin {
+ let Spellings = ["__builtin_spirv_workitem_id_y"];
+ let Attributes = [NoThrow, Const];
+ let Prototype = "void(...)";
+}
+
+def SPIRVGetWrkitemIdZ : Builtin {
+ let Spellings = ["__builtin_spirv_workitem_id_z"];
+ let Attributes = [NoThrow, Const];
+ let Prototype = "void(...)";
+}
+
+def SPIRVBallot : Builtin {
+ let Spellings = ["__builtin_spirv_ballot"];
+ let Attributes = [NoThrow, Const];
+ let Prototype = "void(...)";
+}
+
+def SPIRVSyncThreads : Builtin {
+ let Spellings = ["__builtin_spirv_sync_threads"];
+ let Attributes = [NoThrow, Const];
+ let Prototype = "void(...)";
+}
diff --git a/clang/lib/Headers/CMakeLists.txt b/clang/lib/Headers/CMakeLists.txt
index acf49e40c447e..fe2f3e6305bff 100644
--- a/clang/lib/Headers/CMakeLists.txt
+++ b/clang/lib/Headers/CMakeLists.txt
@@ -289,6 +289,7 @@ set(gpu_files
gpuintrin.h
nvptxintrin.h
amdgpuintrin.h
+ spirvintrin.h
)
set(windows_only_files
diff --git a/clang/lib/Headers/gpuintrin.h b/clang/lib/Headers/gpuintrin.h
index 0fb3916acac61..9f61ff5095824 100644
--- a/clang/lib/Headers/gpuintrin.h
+++ b/clang/lib/Headers/gpuintrin.h
@@ -60,6 +60,8 @@ _Pragma("omp end declare target");
#include <nvptxintrin.h>
#elif defined(__AMDGPU__)
#include <amdgpuintrin.h>
+#elif defined(__SPIRV__)
+#include <spirvintrin.h>
#elif !defined(_OPENMP)
#error "This header is only meant to be used on GPU architectures."
#endif
diff --git a/clang/lib/Headers/spirvintrin.h b/clang/lib/Headers/spirvintrin.h
new file mode 100644
index 0000000000000..5c13a6a592ce0
--- /dev/null
+++ b/clang/lib/Headers/spirvintrin.h
@@ -0,0 +1,186 @@
+//===-- spirvintrin.h - SPIR-V intrinsic functions ------------------------===//
+//
+// Part of the LLVM Project, under the Apache 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 __SPIRVINTRIN_H
+#define __SPIRVINTRIN_H
+
+#ifndef __SPIRV__
+#error "This file is intended for SPIRV targets"
+#endif
+
+#ifndef __GPUINTRIN_H
+#error "Never use <spirvintrin.h> directly; include <gpuintrin.h> instead"
+#endif
+
+_Pragma("omp begin declare target device_type(nohost)");
+_Pragma("omp begin declare variant match(device = {arch(amdgcn)})");
+
+// Type aliases to the address spaces used by the SPIRV backend.
+#define __gpu_private __attribute__((address_space(5)))
+#define __gpu_constant __attribute__((address_space(4)))
+#define __gpu_local __attribute__((address_space(3)))
+#define __gpu_global __attribute__((address_space(1)))
+#define __gpu_generic __attribute__((address_space(0)))
+
+// Attribute to declare a function as a kernel.
+#define __gpu_kernel __attribute__((spir_kernel, visibility("protected")))
+
+// Returns the number of workgroups in the 'x' dimension of the grid.
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_blocks_x(void) {
+ return __builtin_spirv_get_num_workgroups_x();
+}
+
+// Returns the number of workgroups in the 'y' dimension of the grid.
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_blocks_y(void) {
+ return __builtin_spirv_get_num_workgroups_y();
+}
+
+// Returns the number of workgroups in the 'z' dimension of the grid.
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_blocks_z(void) {
+ return __builtin_spirv_get_num_workgroups_z();
+}
+
+// Returns the 'x' dimension of the current SPIR-V workgroup's id.
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_block_id_x(void) {
+ return __builtin_spirv_get_workgroup_id_x();
+}
+
+// Returns the 'y' dimension of the current SPIR-V workgroup's id.
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_block_id_y(void) {
+ return __builtin_spirv_get_workgroup_id_y();
+}
+
+// Returns the 'z' dimension of the current SPIR-V workgroup's id.
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_block_id_z(void) {
+ return __builtin_spirv_get_workgroup_id_z();
+}
+
+// Returns the number of workitems in the 'x' dimension.
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_threads_x(void) {
+ return __builtin_spirv_workgroup_size_x();
+}
+
+// Returns the number of workitems in the 'y' dimension.
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_threads_y(void) {
+ return __builtin_spirv_workgroup_size_y();
+}
+
+// Returns the number of workitems in the 'z' dimension.
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_threads_z(void) {
+ return __builtin_spirv_workgroup_size_z();
+}
+
+// Returns the 'x' dimension id of the workitem in the current SPIR-V workgroup.
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_thread_id_x(void) {
+ return __builtin_spirv_workitem_id_x();
+}
+
+// Returns the 'y' dimension id of the workitem in the current SPIR-V workgroup.
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_thread_id_y(void) {
+ return __builtin_spirv_workitem_id_y();
+}
+
+// Returns the 'z' dimension id of the workitem in the current SPIR-V workgroup.
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_thread_id_z(void) {
+ return __builtin_spirv_workitem_id_z();
+}
+
+// Returns the size of a wavefront
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_lanes(void) {
+ __builtin_unreachable();
+}
+
+// Returns the id of the thread inside of a wavefront executing together.
+_DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_lane_id(void) {
+ __builtin_unreachable();
+}
+
+// Returns the bit-mask of active threads in the current wavefront.
+_DEFAULT_FN_ATTRS static __inline__ uint64_t __gpu_lane_mask(void) {
+ __builtin_unreachable();
+}
+
+// Copies the value from the first active thread in the wavefront to the rest.
+_DEFAULT_FN_ATTRS static __inline__ uint32_t
+__gpu_read_first_lane_u32(uint64_t __lane_mask, uint32_t __x) {
+ __builtin_unreachable();
+}
+
+// Returns a bitmask of threads in the current lane for which \p x is true.
+_DEFAULT_FN_ATTRS static __inline__ uint64_t __gpu_ballot(uint64_t __lane_mask,
+ bool __x) {
+ return __lane_mask & __builtin_spirv_ballot(__x);
+}
+
+// Waits for all the threads in the block to converge and issues a fence.
+_DEFAULT_FN_ATTRS static __inline__ void __gpu_sync_threads(void) {
+ __builtin_spirv_sync_threads();
+}
+
+// Wait for all threads in the wavefront to converge.
+_DEFAULT_FN_ATTRS static __inline__ void __gpu_sync_lane(uint64_t __lane_mask) {
+ __builtin_unreachable();
+}
+
+// Shuffles the the lanes inside the wavefront according to the given index.
+_DEFAULT_FN_ATTRS static __inline__ uint32_t
+__gpu_shuffle_idx_u32(uint64_t __lane_mask, uint32_t __idx, uint32_t __x,
+ uint32_t __width) {
+ __builtin_unreachable();
+}
+
+// Returns a bitmask marking all lanes that have the same value of __x.
+_DEFAULT_FN_ATTRS static __inline__ uint64_t
+__gpu_match_any_u32(uint64_t __lane_mask, uint32_t __x) {
+ return __gpu_match_any_u32_impl(__lane_mask, __x);
+}
+
+// Returns a bitmask marking all lanes that have the same value of __x.
+_DEFAULT_FN_ATTRS static __inline__ uint64_t
+__gpu_match_any_u64(uint64_t __lane_mask, uint64_t __x) {
+ return __gpu_match_any_u64_impl(__lane_mask, __x);
+}
+
+// Returns the current lane mask if every lane contains __x.
+_DEFAULT_FN_ATTRS static __inline__ uint64_t
+__gpu_match_all_u32(uint64_t __lane_mask, uint32_t __x) {
+ return __gpu_match_all_u32_impl(__lane_mask, __x);
+}
+
+// Returns the current lane mask if every lane contains __x.
+_DEFAULT_FN_ATTRS static __inline__ uint64_t
+__gpu_match_all_u64(uint64_t __lane_mask, uint64_t __x) {
+ return __gpu_match_all_u64_impl(__lane_mask, __x);
+}
+
+// Returns true if the flat pointer points to SPIRV 'shared' memory.
+_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_local(void *ptr) {
+ return __builtin_spirv_is_shared((void [[clang::address_space(0)]] *)((
+ void [[clang::opencl_generic]] *)ptr));
+}
+
+// Returns true if the flat pointer points to SPIRV 'private' memory.
+_DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_private(void *ptr) {
+ return __builtin_spirv_is_private((void [[clang::address_space(0)]] *)((
+ void [[clang::opencl_generic]] *)ptr));
+}
+
+// Terminates execution of the associated wavefront.
+_DEFAULT_FN_ATTRS [[noreturn]] static __inline__ void __gpu_exit(void) {
+ __builtin_unreachable();
+}
+
+// Suspend the thread briefly to assist the scheduler during busy loops.
+_DEFAULT_FN_ATTRS static __inline__ void __gpu_thread_suspend(void) {
+ // no op
+}
+
+_Pragma("omp end declare variant");
+_Pragma("omp end declare target");
+
+#endif // __SPIRVINTRIN_H
diff --git a/libc/benchmarks/gpu/timing/spirv64/CMakeLists.txt b/libc/benchmarks/gpu/timing/spirv64/CMakeLists.txt
new file mode 100644
index 0000000000000..506801b115f23
--- /dev/null
+++ b/libc/benchmarks/gpu/timing/spirv64/CMakeLists.txt
@@ -0,0 +1,11 @@
+add_header_library(
+ spirv64_timing
+ HDRS
+ timing.h
+ DEPENDS
+ libc.src.__support.common
+ libc.src.__support.macros.config
+ libc.src.__support.macros.attributes
+ libc.src.__support.CPP.type_traits
+ libc.src.__support.CPP.array
+)
diff --git a/libc/benchmarks/gpu/timing/spirv64/timing.h b/libc/benchmarks/gpu/timing/spirv64/timing.h
new file mode 100644
index 0000000000000..d8b4ce19d3814
--- /dev/null
+++ b/libc/benchmarks/gpu/timing/spirv64/timing.h
@@ -0,0 +1,152 @@
+//===------------- SPIR-V implementation of timing utils --------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache 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_UTILS_GPU_TIMING_SPIRV64
+#define LLVM_LIBC_UTILS_GPU_TIMING_SPIRV64
+
+#include "src/__support/CPP/array.h"
+#include "src/__support/CPP/type_traits.h"
+#include "src/__support/GPU/utils.h"
+#include "src/__support/common.h"
+#include "src/__support/macros/attributes.h"
+#include "src/__support/macros/config.h"
+
+#include <stdint.h>
+
+namespace LIBC_NAMESPACE_DECL {
+
+// Returns the overhead associated with calling the profiling region. This
+// allows us to substract the constant-time overhead from the latency to
+// obtain a true result. This can vary with system load.
+[[gnu::noinline]] static uint64_t overhead() {
+ volatile uint32_t x = 1;
+ uint32_t y = x;
+ uint64_t start = gpu::processor_clock();
+ asm("" ::"llr"(start));
+ uint32_t result = y;
+ asm("or.b32 %[v_reg], %[v_reg], 0;" ::[v_reg] "r"(result));
+ uint64_t stop = gpu::processor_clock();
+ volatile auto storage = result;
+ return stop - start;
+}
+
+// Stimulate a simple function and obtain its latency in clock cycles on the
+// system. This function cannot be inlined or else it will disturb the very
+// delicate balance of hard-coded dependencies.
+template <typename F, typename T>
+[[gnu::noinline]] static LIBC_INLINE uint64_t latency(F f, T t) {
+ // We need to store the input somewhere to guarantee that the compiler will
+ // not constant propagate it and remove the profiling region.
+ volatile T storage = t;
+ T arg = storage;
+
+ // Get the current timestamp from the clock.
+ gpu::memory_fence();
+ uint64_t start = gpu::processor_clock();
+
+ // This forces the compiler to load the input argument and run the clock cycle
+ // counter before the profiling region.
+ asm("" ::"llr"(start));
+
+ // Run the function under test and return its value.
+ auto result = f(arg);
+
+ // This inline assembly performs a no-op which forces the result to both be
+ // used and prevents us from exiting this region before it's complete.
+ asm("or.b32 %[v_reg], %[v_reg], 0;" ::[v_reg] "r"(result));
+
+ // Obtain the current timestamp after running the calculation and force
+ // ordering.
+ uint64_t stop = gpu::processor_clock();
+ gpu::memory_fence();
+ asm("" ::"r"(stop));
+ volatile T output = result;
+
+ // Return the time elapsed.
+ return stop - start;
+}
+
+template <typename F, typename T1, typename T2>
+static LIBC_INLINE uint64_t latency(F f, T1 t1, T2 t2) {
+ volatile T1 storage = t1;
+ volatile T2 storage2 = t2;
+ T1 arg = storage;
+ T2 arg2 = storage2;
+
+ gpu::memory_fence();
+ uint64_t start = gpu::processor_clock();
+
+ asm("" ::"llr"(start));
+
+ auto result = f(arg, arg2);
+
+ asm("or.b32 %[v_reg], %[v_reg], 0;" ::[v_reg] "r"(result));
+
+ uint64_t stop = gpu::processor_clock();
+ gpu::memory_fence();
+ asm("" ::"r"(stop));
+ volatile auto output = result;
+
+ return stop - start;
+}
+
+// Provides throughput benchmarking.
+template <typename F, typename T, size_t N>
+[[gnu::noinline]] static LIBC_INLINE uint64_t
+throughput(F f, const cpp::array<T, N> &inputs) {
+ asm("" ::"r"(&inputs));
+
+ gpu::memory_fence();
+ uint64_t start = gpu::processor_clock();
+
+ asm("" ::"llr"(start));
+
+ uint64_t result;
+ for (auto input : inputs) {
+ asm("" ::"r"(input));
+ result = f(input);
+ asm("" ::"r"(result));
+ }
+
+ uint64_t stop = gpu::processor_clock();
+ gpu::memory_fence();
+ asm("" ::"r"(stop));
+ volatile auto output = result;
+
+ // Return the time elapsed.
+ return stop - start;
+}
+
+// Provides throughput benchmarking for 2 arguments (e.g. atan2())
+template <typename F, typename T, size_t N>
+[[gnu::noinline]] static LIBC_INLINE uint64_t throughput(
+ F f, const cpp::array<T, N> &inputs1, const cpp::array<T, N> &inputs2) {
+ asm("" ::"r"(&inputs1), "r"(&inputs2));
+
+ gpu::memory_fence();
+ uint64_t start = gpu::processor_clock();
+
+ asm("" ::"llr"(start));
+
+ uint64_t result;
+ for (size_t i = 0; i < inputs1.size(); i++) {
+ result = f(inputs1[i], inputs2[i]);
+ asm("" ::"r"(result));
+ }
+
+ uint64_t stop = gpu::processor_clock();
+ gpu::memory_fence();
+ asm("" ::"r"(stop));
+ volatile auto output = result;
+
+ // Return the time elapsed.
+ return stop - start;
+}
+} // namespace LIBC_NAMESPACE_DECL
+
+#endif // LLVM_LIBC_UTILS_GPU_TIMING_SPIRV64
diff --git a/libc/cmake/modules/LLVMLibCArchitectures.cmake b/libc/cmake/modules/LLVMLibCArchitectures.cmake
index a5286ee9e96db..31c10e8433cb3 100644
--- a/libc/cmake/modules/LLVMLibCArchitectures.cmake
+++ b/libc/cmake/modules/LLVMLibCArchitectures.cmake
@@ -73,8 +73,9 @@ function(get_arch_and_system_from_triple triple arch_var sys_var)
# Setting OS name for GPU architectures.
list(GET triple_comps -1 gpu_target_sys)
+ list(GET triple_comps 1 vendor)
if(gpu_target_sys MATCHES "^amdhsa" OR gpu_target_sys MATCHES "^cuda" OR
- gpu_target_sys MATCHES "^opencl")
+ (vendor STREQUAL "intel" AND gpu_target_sys STREQUAL "unknown"))
set(target_sys "gpu")
endif()
diff --git a/libc/config/gpu/spirv64/config.json b/libc/config/gpu/spirv64/config.json
new file mode 100644
index 0000000000000..30ae10e2cfd61
--- /dev/null
+++ b/libc/config/gpu/spirv64/config.json
@@ -0,0 +1,40 @@
+{
+ "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
+ },
+ "LIBC_CONF_PRINTF_RUNTIME_DISPATCH": {
+ "value": false
+ }
+ },
+ "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..15c4b51cb2410
--- /dev/null
+++ b/libc/config/gpu/spirv64/entrypoints.txt
@@ -0,0 +1,610 @@
+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.memalignment
+ 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
+ libc.src.time.mktime
+
+ # 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.fmaf16
+ 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..69bb5e1e1c473
--- /dev/null
+++ b/libc/src/math/spirv64/CMakeLists.txt
@@ -0,0 +1,482 @@
+# Math functions not yet available in the libc project, or those not yet tuned
+# for GPU workloads are provided as wrappers over vendor libraries. If we find
+# them ahead of time we will import them statically. Otherwise, we will keep
+# them as external references and expect them to be resolved by the user when
+# they compile. In the future,we will use implementations from the 'libc'
+# project and not provide these wrappers.
+
+add_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
+)
+
+# The following functions currently are not implemented natively and borrow from
+# existing implementations. This will be removed in the future.
+add_entrypoint_object(
+ acos
+ SRCS
+ acos.cpp
+ HDRS
+ ../acos.h
+ COMPILE_OPTIONS
+ ${bitcode_link_flags}
+ -O2
+ VENDOR
+)
+
+add_entrypoint_object(
+ 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(
+ cos
+ SRCS
+ cos.cpp
+ HDRS
+ ../cos.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..78b7159491a1c
--- /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 __builtin_acos(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..73e26f563b788
--- /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 __builtin_acosf(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..3821989ff1970
--- /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 __builtin_acosh(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..bbace2af36064
--- /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 __builtin_acoshf(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..16633dc418aaf
--- /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 __builtin_asin(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..c8513e1bce303
--- /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 __builtin_asinf(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..36bb795cccf44
--- /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 __builtin_asinh(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..b3e51a2bf416f
--- /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 __builtin_asinhf(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..00263df7c109b
--- /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 __builtin_atan(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..9701e3b4933ba
--- /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 __builtin_atan2(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..a74174512df07
--- /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 __builtin_atan2f(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..e1bbf53e43e78
--- /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 __builtin_atanf(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..ce601341653e2
--- /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 __builtin_atanh(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..ded8728e09006
--- /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 __builtin_atanhf(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..d91ac1c1de425
--- /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 __builtin_cos(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..d2ed7602169d0
--- /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 __builtin_cosf(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..07929c897d8bb
--- /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 __builtin_cosh(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..c6b20c237372c
--- /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 __builtin_coshf(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..24fa43f2b09f7
--- /dev/null
+++ b/libc/src/math/spirv64/declarations.h
@@ -0,0 +1,21 @@
+//===-- SPIR-V 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_SPIRV64_DECLARATIONS_H
+#define LLVM_LIBC_SRC_MATH_SPIRV64_DECLARATIONS_H
+
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+
+extern "C" {
+}
+
+} // namespace LIBC_NAMESPACE_DECL
+
+#endif // LLVM_LIBC_SRC_MATH_SPIRV64_DECLARATIONS_H
diff --git a/libc/src/math/spirv64/erf.cpp b/libc/src/math/spirv64/erf.cpp
new file mode 100644
index 0000000000000..fd223f681f32b
--- /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 __builtin_erf(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..e9584cf9e254b
--- /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 __builtin_erff(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..2bba5c8fea8c9
--- /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 __builtin_exp(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..bd93686684b37
--- /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 __builtin_exp10(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..4a05a15dcecfe
--- /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 __builtin_exp10f(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..aaad23e0b8167
--- /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 __builtin_exp2(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..98938cd9ef266
--- /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 __builtin_exp2f(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..ef2a7434876d5
--- /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 __builtin_expf(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..a69126ce7eb3f
--- /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 __builtin_expm1(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..461aff44ce3a3
--- /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 __builtin_expm1f(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..bf6b1a2512760
--- /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 __builtin_fdim(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..1c66f4c89c826
--- /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 __builtin_fdimf(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..a2496fecb01c4
--- /dev/null
+++ b/libc/src/math/spirv64/fmax.cpp
@@ -0,0 +1,20 @@
+//===-- Implementation of the fmax function for GPU -----------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/math/fmax.h"
+
+#include "src/__support/common.h"
+#include "src/__support/macros/config.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..586fd32fee59b
--- /dev/null
+++ b/libc/src/math/spirv64/fmaxf.cpp
@@ -0,0 +1,22 @@
+//===-- 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/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(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..02229f8a0e901
--- /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 __builtin_hypot(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..1af9cc669f2b3
--- /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 __builtin_hypotf(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..7bde3c8bd1d01
--- /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 __builtin_ilogb(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..de035db7e1959
--- /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 __builtin_ilogbf(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..d3bb12f76d034
--- /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 __builtin_lgamma(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..8d18e1ac8f0d6
--- /dev/null
+++ b/libc/src/math/spirv64/lgamma_r.cpp
@@ -0,0 +1,23 @@
+//===-- 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)) {
+ double result = __builtin_lgamma(x);
+ *signp = (result < 0.0) ? -1 : 1;
+ return result;
+}
+
+} // 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..62a8731d350fe
--- /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 __builtin_log(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..ebff4cd9fc98d
--- /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 __builtin_log10(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..cc25c1b81a25f
--- /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 __builtin_log10f(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..b95493363ea63
--- /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 __builtin_log1p(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..85fade432d3e3
--- /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 __builtin_log1pf(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..be39ab4807805
--- /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 __builtin_log2(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..f106f610721f9
--- /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 __builtin_log2f(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..d273d4b39b0e2
--- /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 __builtin_logb(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..3679b4dfc1ce3
--- /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 __builtin_logbf(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..41d89b3cdc9ac
--- /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 __builtin_logf(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..e2a6cbda6ab10
--- /dev/null
+++ b/libc/src/math/spirv64/lrintf.cpp
@@ -0,0 +1,19 @@
+//===-- 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 __builtin_lrintf(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..c85b2dd74b73f
--- /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 __builtin_nextafter(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..7d56b4c004f96
--- /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 __builtin_nextafterf(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..68639ef3f4b88
--- /dev/null
+++ b/libc/src/math/spirv64/platform.h
@@ -0,0 +1,103 @@
+//===-- SPIR-V specific definitions for math support ----------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_LIBC_SRC_MATH_SPIRV64_SPIRV_H
+#define LLVM_LIBC_SRC_MATH_SPIRV64_SPIRV_H
+
+#include "declarations.h"
+
+#include "src/__support/macros/attributes.h"
+#include "src/__support/macros/config.h"
+
+namespace LIBC_NAMESPACE_DECL {
+namespace internal {
+LIBC_INLINE double acos(double x) { return __builtin_acos(x); }
+LIBC_INLINE float acosf(float x) { return __builtin_acosf(x); }
+LIBC_INLINE double acosh(double x) { return __builtin_acosh(x); }
+LIBC_INLINE float acoshf(float x) { return __builtin_acoshf(x); }
+LIBC_INLINE double asin(double x) { return __builtin_asin(x); }
+LIBC_INLINE float asinf(float x) { return __builtin_asinf(x); }
+LIBC_INLINE double asinh(double x) { return __builtin_asinh(x); }
+LIBC_INLINE float asinhf(float x) { return __builtin_asinhf(x); }
+LIBC_INLINE double atan2(double x, double y) { return __builtin_atan2(x, y); }
+LIBC_INLINE float atan2f(float x, float y) { return __builtin_atan2f(x, y); }
+LIBC_INLINE double atan(double x) { return __builtin_atan(x); }
+LIBC_INLINE float atanf(float x) { return __builtin_atanf(x); }
+LIBC_INLINE double atanh(double x) { return __builtin_atanh(x); }
+LIBC_INLINE float atanhf(float x) { return __builtin_atanhf(x); }
+LIBC_INLINE double cos(double x) { return __builtin_cos(x); }
+LIBC_INLINE float cosf(float x) { return __builtin_cosf(x); }
+LIBC_INLINE double cosh(double x) { return __builtin_cosh(x); }
+LIBC_INLINE float coshf(float x) { return __builtin_coshf(x); }
+LIBC_INLINE double erf(double x) { return __builtin_erf(x); }
+LIBC_INLINE float erff(float x) { return __builtin_erff(x); }
+LIBC_INLINE double exp(double x) { return __builtin_exp(x); }
+LIBC_INLINE float expf(float x) { return __builtin_expf(x); }
+LIBC_INLINE double exp2(double x) { return __builtin_exp2(x); }
+LIBC_INLINE float exp2f(float x) { return __builtin_exp2f(x); }
+LIBC_INLINE double exp10(double x) { return __builtin_exp10(x); }
+LIBC_INLINE float exp10f(float x) { return __builtin_exp10f(x); }
+LIBC_INLINE double expm1(double x) { return __builtin_expm1(x); }
+LIBC_INLINE float expm1f(float x) { return __builtin_expm1f(x); }
+LIBC_INLINE double fdim(double x, double y) { return __builtin_fdim(x, y); }
+LIBC_INLINE float fdimf(float x, float y) { return __builtin_fdimf(x, y); }
+LIBC_INLINE double hypot(double x, double y) { return __builtin_hypot(x, y); }
+LIBC_INLINE float hypotf(float x, float y) { return __builtin_hypotf(x, y); }
+LIBC_INLINE int ilogb(double x) { return __builtin_ilogb(x); }
+LIBC_INLINE int ilogbf(float x) { return __builtin_ilogbf(x); }
+LIBC_INLINE double ldexp(double x, int i) { return __builtin_ldexp(x, i); }
+LIBC_INLINE float ldexpf(float x, int i) { return __builtin_ldexpf(x, i); }
+LIBC_INLINE long long llrint(double x) { return __builtin_llrint(x); }
+LIBC_INLINE long long llrintf(float x) { return __builtin_llrintf(x); }
+LIBC_INLINE double log10(double x) { return __builtin_log10(x); }
+LIBC_INLINE float log10f(float x) { return __builtin_log10f(x); }
+LIBC_INLINE double log1p(double x) { return __builtin_log1p(x); }
+LIBC_INLINE float log1pf(float x) { return __builtin_log1pf(x); }
+LIBC_INLINE double log2(double x) { return __builtin_log2(x); }
+LIBC_INLINE float log2f(float x) { return __builtin_log2f(x); }
+LIBC_INLINE double log(double x) { return __builtin_log(x); }
+LIBC_INLINE float logf(float x) { return __builtin_logf(x); }
+LIBC_INLINE long lrint(double x) { return __builtin_lrint(x); }
+LIBC_INLINE long lrintf(float x) { return __builtin_lrintf(x); }
+LIBC_INLINE double nextafter(double x, double y) {
+ return __builtin_nextafter(x, y);
+}
+LIBC_INLINE float nextafterf(float x, float y) { return __builtin_nextafterf(x, y); }
+LIBC_INLINE double pow(double x, double y) { return __builtin_pow(x, y); }
+LIBC_INLINE float powf(float x, float y) { return __builtin_powf(x, y); }
+LIBC_INLINE double sin(double x) { return __builtin_sin(x); }
+LIBC_INLINE float sinf(float x) { return __builtin_sinf(x); }
+LIBC_INLINE void sincos(double x, double *sinptr, double *cosptr) {
+ return __builtin_sincos(x, sinptr, cosptr);
+}
+LIBC_INLINE void sincosf(float x, float *sinptr, float *cosptr) {
+ return __builtin_sincosf(x, sinptr, cosptr);
+}
+LIBC_INLINE double sinh(double x) { return __builtin_sinh(x); }
+LIBC_INLINE float sinhf(float x) { return __builtin_sinhf(x); }
+LIBC_INLINE double tan(double x) { return __builtin_tan(x); }
+LIBC_INLINE float tanf(float x) { return __builtin_tanf(x); }
+LIBC_INLINE double tanh(double x) { return __builtin_tanh(x); }
+LIBC_INLINE float tanhf(float x) { return __builtin_tanhf(x); }
+LIBC_INLINE double scalbn(double x, int i) { return __builtin_scalbn(x, i); }
+LIBC_INLINE float scalbnf(float x, int i) { return __builtin_scalbnf(x, i); }
+LIBC_INLINE double frexp(double x, int *i) { return __builtin_frexp(x, i); }
+LIBC_INLINE float frexpf(float x, int *i) { return __builtin_frexpf(x, i); }
+LIBC_INLINE double remquo(double x, double y, int *i) {
+ return __builtin_remquo(x, y, i);
+}
+LIBC_INLINE float remquof(float x, float y, int *i) {
+ return __builtin_remquof(x, y, i);
+}
+LIBC_INLINE double tgamma(double x) { return __builtin_tgamma(x); }
+LIBC_INLINE float tgammaf(float x) { return __builtin_tgammaf(x); }
+
+} // namespace internal
+} // namespace LIBC_NAMESPACE_DECL
+
+#endif // LLVM_LIBC_SRC_MATH_SPIRV64_SPIRV_H
diff --git a/libc/src/math/spirv64/powf.cpp b/libc/src/math/spirv64/powf.cpp
new file mode 100644
index 0000000000000..3cc6ab15fb706
--- /dev/null
+++ b/libc/src/math/spirv64/powf.cpp
@@ -0,0 +1,19 @@
+//===-- 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 __builtin_powf(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..a1601743f3738
--- /dev/null
+++ b/libc/src/math/spirv64/powi.cpp
@@ -0,0 +1,19 @@
+//===-- 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 __builtin_powi(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..dfe40b6aa8ac0
--- /dev/null
+++ b/libc/src/math/spirv64/powif.cpp
@@ -0,0 +1,19 @@
+//===-- Implementation of the powif 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 __builtin_powif(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..c2697311edc5c
--- /dev/null
+++ b/libc/src/math/spirv64/remquo.cpp
@@ -0,0 +1,21 @@
+//===-- 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)) {
+ return __builtin_remquo(x, y, quo);
+}
+
+} // 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..a5b4504bad748
--- /dev/null
+++ b/libc/src/math/spirv64/remquof.cpp
@@ -0,0 +1,21 @@
+//===-- 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)) {
+ return __builtin_remquof(x, y, quo);
+}
+
+} // 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..442c474e561ec
--- /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_scalbn(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..68091d2ad4d60
--- /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_scalbnf(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..8c3e408d43c00
--- /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 __builtin_sin(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..37d2252079b6e
--- /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)) {
+ return __builtin_sincos(x, sinptr, 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..77e8eed1efba5
--- /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)) {
+ return __builtin_sincosf(x, sinptr, 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..f27d59ef2b3b2
--- /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 __builtin_sinf(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..92086f10fe8e9
--- /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 __builtin_sinh(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..9198957e70ed4
--- /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 __builtin_sinhf(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..e2bea0c207d77
--- /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 __builtin_tan(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..08e35e381fae2
--- /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 __builtin_tanf(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..62601e82cae09
--- /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 __builtin_tanh(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..2a5c18a2d7f6b
--- /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 __builtin_tanhf(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..16fc5abaab036
--- /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 __builtin_tgamma(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..ddafff1eec664
--- /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 __builtin_tgammaf(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..4c722c0208cba
--- /dev/null
+++ b/libc/startup/gpu/spirv64/start.cpp
@@ -0,0 +1,74 @@
+//===-- Implementation of crt for spirv64 ---------------------------------===//
+//
+// Part of the LLVM Project, under the Apache 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 {
+
+DataEnvironment app;
+
+// FIXME: Factor this out into common logic so we don't need to stub it here.
+void teardown_main_tls() {}
+
+extern "C" {
+uintptr_t *__init_array_start [[gnu::visibility("protected")]];
+uintptr_t *__init_array_end [[gnu::visibility("protected")]];
+uintptr_t *__fini_array_start [[gnu::visibility("protected")]];
+uintptr_t *__fini_array_end [[gnu::visibility("protected")]];
+}
+
+using InitCallback = void(void);
+using FiniCallback = void(void);
+
+static void call_init_array_callbacks(int, char **, char **) {
+ 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])();
+}
+
+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
+
+extern "C" [[gnu::visibility("protected"), clang::spir_kernel]] 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::spir_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::spir_kernel]] void
+_end(int retval) {
+ // To finis the execution we invoke all the callbacks registered via 'atexit'
+ // and then exit with the appropriate return value.
+ LIBC_NAMESPACE::exit(retval);
+}
More information about the cfe-commits
mailing list