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