From 2d9c400e16778a26c6034cd2e2a6b1773a282ce8 Mon Sep 17 00:00:00 2001 From: Jon Chesterfield 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 *); +double __ocml_remquo_f64(double, double, gpu::Private *); +double __ocml_tgamma_f64(double); +float __ocml_tgamma_f32(float); +double __ocml_lgamma_f64(double); +double __ocml_lgamma_r_f64(double, gpu::Private *); +} + +} // 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 *)&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(__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(__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(__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(__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 + +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 *)&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 *)&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(__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(__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(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 &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 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( \ @@ -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 __omp_rtl_device_environment; + gnu::visibility("protected")]] Constant __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 Data[state::SharedScratchpadSize]; [[gnu::aligned( - allocator::ALIGNMENT)]] unsigned char Usage[mapping::MaxThreadsPerTeam]; + allocator::ALIGNMENT)]] + [[clang::loader_uninitialized]] + static + Local Usage[mapping::MaxThreadsPerTeam]; }; +Local SharedMemorySmartStackTy::Data[state::SharedScratchpadSize]; +Local 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 - 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( memory::allocGlobal(sizeof(ThreadStateTy), "ThreadStates alloc")); - uintptr_t *ThreadStatesBitsPtr = reinterpret_cast(&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; }