[llvm-branch-commits] [clang] [llvm] [Offload][Lang] Add math function declarations (PR #217760)
Sophia Herrmann via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Aug 20 17:53:57 PDT 2026
https://github.com/jellytabby updated https://github.com/llvm/llvm-project/pull/217760
>From ef55c1e57c5b8bd8bbf6a72e150688bc485b29b2 Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Wed, 19 Aug 2026 14:40:52 -0700
Subject: [PATCH] add math functions
---
clang/lib/Driver/ToolChains/Clang.cpp | 107 +++++---
.../test/CodeGenCUDA/offload_via_llvm_math.cu | 252 ++++++++++++++++++
.../lib/amdgpu-amd-amdhsa/.keep | 1 +
clang/test/Driver/gpu-libc.c | 12 +
offload/languages/include/cuda/cuda_runtime.h | 3 +
offload/languages/include/hip/hip_runtime.h | 3 +
.../languages/include/kernel/LanguageMath.h | 137 ++++++++++
offload/languages/kernel/CMakeLists.txt | 1 +
offload/test/offloading/language/math.cpp | 107 ++++++++
9 files changed, 587 insertions(+), 36 deletions(-)
create mode 100644 clang/test/CodeGenCUDA/offload_via_llvm_math.cu
create mode 100644 clang/test/Driver/Inputs/basic_gpu_tree/lib/amdgpu-amd-amdhsa/.keep
create mode 100644 offload/languages/include/kernel/LanguageMath.h
create mode 100644 offload/test/offloading/language/math.cpp
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 8b21400ab959f..5d898d04d34a7 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -9911,6 +9911,60 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
ShouldForwardForToolChain(A, TC);
};
+ bool ShouldLinkGPULibC =
+ Args.hasFlag(options::OPT_offloadlib, OPT_no_offloadlib, true) &&
+ !Args.hasArg(options::OPT_nostdlib, options::OPT_r,
+ options::OPT_nodefaultlibs, options::OPT_nolibc,
+ options::OPT_nogpulibc);
+
+ auto GetGPULibCPathForToolChain =
+ [&](const ToolChain &TC) -> std::optional<std::string> {
+ StringRef TripleString = TC.getTripleString();
+ llvm::Triple Triple(TripleString.ends_with("-llvm")
+ ? TripleString.drop_back(5)
+ : TripleString);
+
+ // GPU libc is installed under the canonical target triple, not the
+ // -foffload-via-llvm internal triple with an LLVM environment suffix.
+ if (Triple.getEnvironment() == llvm::Triple::LLVM)
+ Triple.setEnvironment(llvm::Triple::UnknownEnvironment);
+ else if (std::optional<std::string> Path = TC.getStdlibPath())
+ return Path;
+
+ auto GetPathForTriple =
+ [&](const llvm::Triple &LibCTriple) -> std::optional<std::string> {
+ SmallString<128> Path(C.getDriver().Dir);
+ llvm::sys::path::append(Path, "..", "lib", LibCTriple.str());
+ if (TC.getVFS().exists(Path))
+ return std::string(Path);
+ return std::nullopt;
+ };
+
+ if (std::optional<std::string> Path = GetPathForTriple(Triple))
+ return Path;
+
+ // Handle the current AMDGPU spelling as well as the legacy amdgcn one.
+ if (Triple.getArchName() == "amdgcn") {
+ llvm::Triple Canon(Triple);
+ Canon.setArchName("amdgpu");
+ if (std::optional<std::string> Path = GetPathForTriple(Canon))
+ return Path;
+ }
+
+ return std::nullopt;
+ };
+ auto AddGPULibCLinkerArgs = [&](const ToolChain &TC,
+ ArgStringList &LinkerArgs) {
+ if (!ShouldLinkGPULibC ||
+ (!TC.getTriple().isNVPTX() && !TC.getTriple().isAMDGPU()))
+ return;
+ if (std::optional<std::string> LibCPath = GetGPULibCPathForToolChain(TC)) {
+ LinkerArgs.emplace_back(Args.MakeArgString(Twine("-L") + *LibCPath));
+ LinkerArgs.emplace_back("-lc");
+ LinkerArgs.emplace_back("-lm");
+ }
+ };
+
ArgStringList CmdArgs;
for (Action::OffloadKind Kind : {Action::OFK_Cuda, Action::OFK_OpenMP,
Action::OFK_HIP, Action::OFK_SYCL}) {
@@ -9946,6 +10000,23 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
(TC->getTriple().isAMDGPU() || TC->getTriple().isNVPTX()))
LinkerArgs.emplace_back("-lompdevice");
+ if (Kind == Action::OFK_OpenMP ||
+ TC->getTripleString().ends_with("-llvm"))
+ AddGPULibCLinkerArgs(*TC, LinkerArgs);
+
+ if (Kind == Action::OFK_OpenMP &&
+ (TC->getTriple().isAMDGPU() || TC->getTriple().isNVPTX()) &&
+ ShouldLinkGPULibC) {
+ if (ToolChainHasRT(*TC, "builtins"))
+ LinkerArgs.emplace_back("-lclang_rt.builtins");
+
+ bool HasFlangRT = getToolChain().getVFS().exists(
+ TC->getCompilerRT(Args, "runtime", ToolChain::FT_Static,
+ /*IsFortran=*/true));
+ if (HasFlangRT && C.getDriver().IsFlangMode())
+ LinkerArgs.emplace_back("-lflang_rt.runtime");
+ }
+
// For SPIR-V, pass some extra flags to `spirv-link`, the out-of-tree
// SPIR-V linker. `spirv-link` isn't called in LTO mode so restrict these
// flags to normal compilation.
@@ -10067,42 +10138,6 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
Args.hasArg(options::OPT_save_temps))
CmdArgs.push_back("--save-temps");
- // Pass in the C library for GPUs if present and not disabled.
- if (Args.hasFlag(options::OPT_offloadlib, OPT_no_offloadlib, true) &&
- !Args.hasArg(options::OPT_nostdlib, options::OPT_r,
- options::OPT_nodefaultlibs, options::OPT_nolibc,
- options::OPT_nogpulibc)) {
- forAllAssociatedToolChains(C, JA, getToolChain(), [&](const ToolChain &TC) {
- // The device C library is only available for NVPTX and AMDGPU targets
- // and we only link it by default for OpenMP currently.
- if ((!TC.getTriple().isNVPTX() && !TC.getTriple().isAMDGPU()) ||
- !JA.isHostOffloading(Action::OFK_OpenMP))
- return;
- bool HasLibC = TC.getStdlibIncludePath().has_value();
- if (HasLibC) {
- CmdArgs.push_back(Args.MakeArgString(
- "--device-linker=" + TC.getTripleString() + "=" + "-lc"));
- CmdArgs.push_back(Args.MakeArgString(
- "--device-linker=" + TC.getTripleString() + "=" + "-lm"));
- }
- auto HasCompilerRT = getToolChain().getVFS().exists(
- TC.getCompilerRT(Args, "builtins", ToolChain::FT_Static,
- /*IsFortran=*/false));
- if (HasCompilerRT)
- CmdArgs.push_back(
- Args.MakeArgString("--device-linker=" + TC.getTripleString() + "=" +
- "-lclang_rt.builtins"));
-
- bool HasFlangRT = getToolChain().getVFS().exists(
- TC.getCompilerRT(Args, "runtime", ToolChain::FT_Static,
- /*IsFortran=*/true));
- if (HasFlangRT && C.getDriver().IsFlangMode())
- CmdArgs.push_back(
- Args.MakeArgString("--device-linker=" + TC.getTripleString() + "=" +
- "-lflang_rt.runtime"));
- });
- }
-
// Add the linker arguments to be forwarded by the wrapper.
CmdArgs.push_back(Args.MakeArgString(Twine("--linker-path=") +
LinkCommand->getExecutable()));
diff --git a/clang/test/CodeGenCUDA/offload_via_llvm_math.cu b/clang/test/CodeGenCUDA/offload_via_llvm_math.cu
new file mode 100644
index 0000000000000..cd78276e007b2
--- /dev/null
+++ b/clang/test/CodeGenCUDA/offload_via_llvm_math.cu
@@ -0,0 +1,252 @@
+// RUN: %clangxx --cuda-device-only -S -emit-llvm -foffload-via-llvm \
+// RUN: --offload-arch=sm_90 -nocudalib \
+// RUN: -I %S/../../../offload/languages/include \
+// RUN: -I %S/../../../offload/languages/include/cuda \
+// RUN: -I %S/../../../offload/languages/include/hip \
+// RUN: -I %S/../../../offload/languages/kernel/include \
+// RUN: %s -o - | FileCheck %s
+
+#include "cuda_runtime.h"
+
+__global__ void math_kernel(double *D, float *F, int *I, long *L,
+ long long *LL) {
+ D[0] = acos(D[0]);
+ D[1] = acosh(D[1]);
+ D[2] = asin(D[2]);
+ D[3] = asinh(D[3]);
+ D[4] = atan(D[4]);
+ D[5] = atan2(D[5], D[6]);
+ D[6] = atanh(D[6]);
+ D[7] = cbrt(D[7]);
+ D[8] = ceil(D[8]);
+ D[9] = copysign(D[9], D[10]);
+ D[10] = cos(D[10]);
+ D[11] = cosh(D[11]);
+ D[12] = erf(D[12]);
+ D[13] = erfc(D[13]);
+ D[14] = exp(D[14]);
+ D[15] = exp2(D[15]);
+ D[16] = exp10(D[16]);
+ D[17] = expm1(D[17]);
+ D[18] = fabs(D[18]);
+ D[19] = fdim(D[19], D[20]);
+ D[20] = floor(D[20]);
+ D[21] = fma(D[21], D[22], D[23]);
+ D[22] = fmax(D[22], D[23]);
+ D[23] = fmin(D[23], D[24]);
+ D[24] = fmod(D[24], D[25]);
+ D[25] = frexp(D[25], I);
+ D[26] = hypot(D[26], D[27]);
+ I[1] = ilogb(D[27]);
+ D[28] = ldexp(D[28], I[1]);
+ D[29] = lgamma(D[29]);
+ LL[0] = llrint(D[30]);
+ LL[1] = llround(D[31]);
+ D[30] = log(D[30]);
+ D[31] = log10(D[31]);
+ D[32] = log1p(D[32]);
+ D[33] = log2(D[33]);
+ D[34] = logb(D[34]);
+ L[0] = lrint(D[35]);
+ L[1] = lround(D[36]);
+ D[35] = modf(D[35], &D[36]);
+ D[36] = nearbyint(D[36]);
+ D[37] = nextafter(D[37], D[38]);
+ D[38] = pow(D[38], D[39]);
+ D[39] = remainder(D[39], D[40]);
+ D[40] = remquo(D[40], D[41], &I[2]);
+ D[41] = rint(D[41]);
+ D[42] = round(D[42]);
+ D[43] = roundeven(D[43]);
+ D[44] = scalbln(D[44], L[2]);
+ D[45] = scalbn(D[45], I[3]);
+ D[46] = sin(D[46]);
+ sincos(D[47], &D[47], &D[48]);
+ sincospi(D[49], &D[49], &D[50]);
+ D[51] = sinh(D[51]);
+ D[52] = sqrt(D[52]);
+ D[53] = tan(D[53]);
+ D[54] = tanh(D[54]);
+ D[55] = tgamma(D[55]);
+ D[56] = trunc(D[56]);
+
+ F[0] = acosf(F[0]);
+ F[1] = acoshf(F[1]);
+ F[2] = asinf(F[2]);
+ F[3] = asinhf(F[3]);
+ F[4] = atanf(F[4]);
+ F[5] = atan2f(F[5], F[6]);
+ F[6] = atanhf(F[6]);
+ F[7] = cbrtf(F[7]);
+ F[8] = ceilf(F[8]);
+ F[9] = copysignf(F[9], F[10]);
+ F[10] = cosf(F[10]);
+ F[11] = coshf(F[11]);
+ F[12] = erff(F[12]);
+ F[13] = erfcf(F[13]);
+ F[14] = expf(F[14]);
+ F[15] = exp2f(F[15]);
+ F[16] = exp10f(F[16]);
+ F[17] = expm1f(F[17]);
+ F[18] = fabsf(F[18]);
+ F[19] = fdimf(F[19], F[20]);
+ F[20] = floorf(F[20]);
+ F[21] = fmaf(F[21], F[22], F[23]);
+ F[22] = fmaxf(F[22], F[23]);
+ F[23] = fminf(F[23], F[24]);
+ F[24] = fmodf(F[24], F[25]);
+ F[25] = frexpf(F[25], &I[4]);
+ F[26] = hypotf(F[26], F[27]);
+ I[5] = ilogbf(F[27]);
+ F[28] = ldexpf(F[28], I[5]);
+ F[29] = lgammaf(F[29]);
+ LL[2] = llrintf(F[30]);
+ LL[3] = llroundf(F[31]);
+ F[30] = logf(F[30]);
+ F[31] = log10f(F[31]);
+ F[32] = log1pf(F[32]);
+ F[33] = log2f(F[33]);
+ F[34] = logbf(F[34]);
+ L[3] = lrintf(F[35]);
+ L[4] = lroundf(F[36]);
+ F[35] = modff(F[35], &F[36]);
+ F[36] = nearbyintf(F[36]);
+ F[37] = nextafterf(F[37], F[38]);
+ F[38] = powf(F[38], F[39]);
+ F[39] = remainderf(F[39], F[40]);
+ F[40] = remquof(F[40], F[41], &I[6]);
+ F[41] = rintf(F[41]);
+ F[42] = roundf(F[42]);
+ F[43] = roundevenf(F[43]);
+ F[44] = scalblnf(F[44], L[5]);
+ F[45] = scalbnf(F[45], I[7]);
+ F[46] = sinf(F[46]);
+ sincosf(F[47], &F[47], &F[48]);
+ sincospif(F[49], &F[49], &F[50]);
+ F[51] = sinhf(F[51]);
+ F[52] = sqrtf(F[52]);
+ F[53] = tanf(F[53]);
+ F[54] = tanhf(F[54]);
+ F[55] = tgammaf(F[55]);
+ F[56] = truncf(F[56]);
+}
+
+// CHECK-NOT: __nv_
+// CHECK: call {{.*}} @acos(
+// CHECK: call {{.*}} @acosh(
+// CHECK: call {{.*}} @asin(
+// CHECK: call {{.*}} @asinh(
+// CHECK: call {{.*}} @atan(
+// CHECK: call {{.*}} @atan2(
+// CHECK: call {{.*}} @atanh(
+// CHECK: call {{.*}} @cbrt(
+// CHECK: call {{.*}} @ceil(
+// CHECK: call {{.*}} @copysign(
+// CHECK: call {{.*}} @cos(
+// CHECK: call {{.*}} @cosh(
+// CHECK: call {{.*}} @erf(
+// CHECK: call {{.*}} @erfc(
+// CHECK: call {{.*}} @exp(
+// CHECK: call {{.*}} @exp2(
+// CHECK: call {{.*}} @exp10(
+// CHECK: call {{.*}} @expm1(
+// CHECK: call {{.*}} @fabs(
+// CHECK: call {{.*}} @fdim(
+// CHECK: call {{.*}} @floor(
+// CHECK: call {{.*}} @fma(
+// CHECK: call {{.*}} @fmax(
+// CHECK: call {{.*}} @fmin(
+// CHECK: call {{.*}} @fmod(
+// CHECK: call {{.*}} @frexp(
+// CHECK: call {{.*}} @hypot(
+// CHECK: call {{.*}} @ilogb(
+// CHECK: call {{.*}} @ldexp(
+// CHECK: call {{.*}} @lgamma(
+// CHECK: call {{.*}} @llrint(
+// CHECK: call {{.*}} @llround(
+// CHECK: call {{.*}} @log(
+// CHECK: call {{.*}} @log10(
+// CHECK: call {{.*}} @log1p(
+// CHECK: call {{.*}} @log2(
+// CHECK: call {{.*}} @logb(
+// CHECK: call {{.*}} @lrint(
+// CHECK: call {{.*}} @lround(
+// CHECK: call {{.*}} @modf(
+// CHECK: call {{.*}} @nearbyint(
+// CHECK: call {{.*}} @nextafter(
+// CHECK: call {{.*}} @pow(
+// CHECK: call {{.*}} @remainder(
+// CHECK: call {{.*}} @remquo(
+// CHECK: call {{.*}} @rint(
+// CHECK: call {{.*}} @round(
+// CHECK: call {{.*}} @roundeven(
+// CHECK: call {{.*}} @scalbln(
+// CHECK: call {{.*}} @scalbn(
+// CHECK: call {{.*}} @sin(
+// CHECK: call {{.*}} @sincos(
+// CHECK: call {{.*}} @sincospi(
+// CHECK: call {{.*}} @sinh(
+// CHECK: call {{.*}} @sqrt(
+// CHECK: call {{.*}} @tan(
+// CHECK: call {{.*}} @tanh(
+// CHECK: call {{.*}} @tgamma(
+// CHECK: call {{.*}} @trunc(
+// CHECK: call {{.*}} @acosf(
+// CHECK: call {{.*}} @acoshf(
+// CHECK: call {{.*}} @asinf(
+// CHECK: call {{.*}} @asinhf(
+// CHECK: call {{.*}} @atanf(
+// CHECK: call {{.*}} @atan2f(
+// CHECK: call {{.*}} @atanhf(
+// CHECK: call {{.*}} @cbrtf(
+// CHECK: call {{.*}} @ceilf(
+// CHECK: call {{.*}} @copysignf(
+// CHECK: call {{.*}} @cosf(
+// CHECK: call {{.*}} @coshf(
+// CHECK: call {{.*}} @erff(
+// CHECK: call {{.*}} @erfcf(
+// CHECK: call {{.*}} @expf(
+// CHECK: call {{.*}} @exp2f(
+// CHECK: call {{.*}} @exp10f(
+// CHECK: call {{.*}} @expm1f(
+// CHECK: call {{.*}} @fabsf(
+// CHECK: call {{.*}} @fdimf(
+// CHECK: call {{.*}} @floorf(
+// CHECK: call {{.*}} @fmaf(
+// CHECK: call {{.*}} @fmaxf(
+// CHECK: call {{.*}} @fminf(
+// CHECK: call {{.*}} @fmodf(
+// CHECK: call {{.*}} @frexpf(
+// CHECK: call {{.*}} @hypotf(
+// CHECK: call {{.*}} @ilogbf(
+// CHECK: call {{.*}} @ldexpf(
+// CHECK: call {{.*}} @lgammaf(
+// CHECK: call {{.*}} @llrintf(
+// CHECK: call {{.*}} @llroundf(
+// CHECK: call {{.*}} @logf(
+// CHECK: call {{.*}} @log10f(
+// CHECK: call {{.*}} @log1pf(
+// CHECK: call {{.*}} @log2f(
+// CHECK: call {{.*}} @logbf(
+// CHECK: call {{.*}} @lrintf(
+// CHECK: call {{.*}} @lroundf(
+// CHECK: call {{.*}} @modff(
+// CHECK: call {{.*}} @nearbyintf(
+// CHECK: call {{.*}} @nextafterf(
+// CHECK: call {{.*}} @powf(
+// CHECK: call {{.*}} @remainderf(
+// CHECK: call {{.*}} @remquof(
+// CHECK: call {{.*}} @rintf(
+// CHECK: call {{.*}} @roundf(
+// CHECK: call {{.*}} @roundevenf(
+// CHECK: call {{.*}} @scalblnf(
+// CHECK: call {{.*}} @scalbnf(
+// CHECK: call {{.*}} @sinf(
+// CHECK: call {{.*}} @sincosf(
+// CHECK: call {{.*}} @sincospif(
+// CHECK: call {{.*}} @sinhf(
+// CHECK: call {{.*}} @sqrtf(
+// CHECK: call {{.*}} @tanf(
+// CHECK: call {{.*}} @tanhf(
+// CHECK: call {{.*}} @tgammaf(
+// CHECK: call {{.*}} @truncf(
diff --git a/clang/test/Driver/Inputs/basic_gpu_tree/lib/amdgpu-amd-amdhsa/.keep b/clang/test/Driver/Inputs/basic_gpu_tree/lib/amdgpu-amd-amdhsa/.keep
new file mode 100644
index 0000000000000..8b137891791fe
--- /dev/null
+++ b/clang/test/Driver/Inputs/basic_gpu_tree/lib/amdgpu-amd-amdhsa/.keep
@@ -0,0 +1 @@
+
diff --git a/clang/test/Driver/gpu-libc.c b/clang/test/Driver/gpu-libc.c
index 88f346f32e0b8..1272a612d31bc 100644
--- a/clang/test/Driver/gpu-libc.c
+++ b/clang/test/Driver/gpu-libc.c
@@ -26,7 +26,19 @@
// RUN: --offload-new-driver --rocm-path=%S/Inputs/rocm --sysroot=%S/Inputs/basic_gpu_tree \
// RUN: -ccc-install-dir %S/Inputs/basic_gpu_tree/bin -x hip %s 2>&1 | FileCheck %s --check-prefix=HIP
// HIP-NOT: "--device-linker=amdgcn-amd-amdhsa=-lc"
+// RUN: %clang -### --target=x86_64-unknown-linux-gnu --offload-arch=gfx908 \
+// RUN: -foffload-via-llvm --sysroot=%S/Inputs/basic_gpu_tree \
+// RUN: -ccc-install-dir %S/Inputs/basic_gpu_tree/bin -x hip %s 2>&1 | FileCheck %s --check-prefix=HIP-VIA-LLVM
+// HIP-VIA-LLVM: clang-linker-wrapper{{.*}}"--device-linker=amdgpu-amd-amdhsa-llvm=-L{{.*}}lib{{/|\\\\}}amdgpu-amd-amdhsa"
+// HIP-VIA-LLVM-SAME: "--device-linker=amdgpu-amd-amdhsa-llvm=-lc"
+// HIP-VIA-LLVM-SAME: "--device-linker=amdgpu-amd-amdhsa-llvm=-lm"
// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fgpu-rdc --offload-arch=sm_52 \
// RUN: --cuda-path=%S/Inputs/CUDA_111/usr/local/cuda --sysroot=%S/Inputs/basic_gpu_tree \
// RUN: -ccc-install-dir %S/Inputs/basic_gpu_tree/bin -x cuda %s 2>&1 | FileCheck %s --check-prefix=CUDA
// CUDA-NOT: "--device-linker=nvptx64-nvidia-cuda=-lc"
+// RUN: %clang -### --target=x86_64-unknown-linux-gnu --offload-arch=sm_52 \
+// RUN: -foffload-via-llvm --sysroot=%S/Inputs/basic_gpu_tree \
+// RUN: -ccc-install-dir %S/Inputs/basic_gpu_tree/bin -x cuda %s 2>&1 | FileCheck %s --check-prefix=CUDA-VIA-LLVM
+// CUDA-VIA-LLVM: clang-linker-wrapper{{.*}}"--device-linker=nvptx64-nvidia-cuda-llvm=-L{{.*}}lib{{/|\\\\}}nvptx64-nvidia-cuda"
+// CUDA-VIA-LLVM-SAME: "--device-linker=nvptx64-nvidia-cuda-llvm=-lc"
+// CUDA-VIA-LLVM-SAME: "--device-linker=nvptx64-nvidia-cuda-llvm=-lm"
diff --git a/offload/languages/include/cuda/cuda_runtime.h b/offload/languages/include/cuda/cuda_runtime.h
index a140ce1a1aca0..ba681383bfd67 100644
--- a/offload/languages/include/cuda/cuda_runtime.h
+++ b/offload/languages/include/cuda/cuda_runtime.h
@@ -17,6 +17,9 @@
#include "../kernel/UndefineLanguageNames.inc"
+// we dont rename the math symbols
+#include "../kernel/LanguageMath.h"
+
#undef LANGUAGE
using cudaDeviceProp = cudaDeviceProp_t;
diff --git a/offload/languages/include/hip/hip_runtime.h b/offload/languages/include/hip/hip_runtime.h
index 5c58f829daa84..f98a843f8490f 100644
--- a/offload/languages/include/hip/hip_runtime.h
+++ b/offload/languages/include/hip/hip_runtime.h
@@ -17,6 +17,9 @@
#include "../kernel/UndefineLanguageNames.inc"
+// we dont rename the math symbols
+#include "../kernel/LanguageMath.h"
+
#undef LANGUAGE
#define hipHostMallocDefault hipHostAllocDefault
diff --git a/offload/languages/include/kernel/LanguageMath.h b/offload/languages/include/kernel/LanguageMath.h
new file mode 100644
index 0000000000000..f7ebec03b5433
--- /dev/null
+++ b/offload/languages/include/kernel/LanguageMath.h
@@ -0,0 +1,137 @@
+//===-- LanguageMath.h - Kernel language math declarations ----------------===//
+//
+// Part of the LLVM Project, under the Apache 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_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_MATH_H
+#define LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_MATH_H
+
+#if defined(__CUDA__) || defined(__HIP__)
+
+extern "C" {
+__device__ float acosf(float);
+__device__ double acos(double);
+__device__ float acoshf(float);
+__device__ double acosh(double);
+__device__ float asinf(float);
+__device__ double asin(double);
+__device__ float asinhf(float);
+__device__ double asinh(double);
+__device__ float atanf(float);
+__device__ double atan(double);
+__device__ float atan2f(float, float);
+__device__ double atan2(double, double);
+__device__ float atanhf(float);
+__device__ double atanh(double);
+__device__ float cbrtf(float);
+__device__ double cbrt(double);
+__device__ float ceilf(float);
+__device__ double ceil(double);
+__device__ float copysignf(float, float);
+__device__ double copysign(double, double);
+__device__ float cosf(float);
+__device__ double cos(double);
+__device__ float coshf(float);
+__device__ double cosh(double);
+__device__ float erff(float);
+__device__ double erf(double);
+__device__ float erfcf(float);
+__device__ double erfc(double);
+__device__ float expf(float);
+__device__ double exp(double);
+__device__ float exp2f(float);
+__device__ double exp2(double);
+__device__ float exp10f(float);
+__device__ double exp10(double);
+__device__ float expm1f(float);
+__device__ double expm1(double);
+__device__ float fabsf(float);
+__device__ double fabs(double);
+__device__ float fdimf(float, float);
+__device__ double fdim(double, double);
+__device__ float floorf(float);
+__device__ double floor(double);
+__device__ float fmaf(float, float, float);
+__device__ double fma(double, double, double);
+__device__ float fmaxf(float, float);
+__device__ double fmax(double, double);
+__device__ float fminf(float, float);
+__device__ double fmin(double, double);
+__device__ float fmodf(float, float);
+__device__ double fmod(double, double);
+__device__ float frexpf(float, int *);
+__device__ double frexp(double, int *);
+__device__ float hypotf(float, float);
+__device__ double hypot(double, double);
+__device__ int ilogbf(float);
+__device__ int ilogb(double);
+__device__ float ldexpf(float, int);
+__device__ double ldexp(double, int);
+__device__ float lgammaf(float);
+__device__ double lgamma(double);
+__device__ long long llrintf(float);
+__device__ long long llrint(double);
+__device__ long long llroundf(float);
+__device__ long long llround(double);
+__device__ float logf(float);
+__device__ double log(double);
+__device__ float log10f(float);
+__device__ double log10(double);
+__device__ float log1pf(float);
+__device__ double log1p(double);
+__device__ float log2f(float);
+__device__ double log2(double);
+__device__ float logbf(float);
+__device__ double logb(double);
+__device__ long lrintf(float);
+__device__ long lrint(double);
+__device__ long lroundf(float);
+__device__ long lround(double);
+__device__ float modff(float, float *);
+__device__ double modf(double, double *);
+__device__ float nearbyintf(float);
+__device__ double nearbyint(double);
+__device__ float nextafterf(float, float);
+__device__ double nextafter(double, double);
+__device__ float powf(float, float);
+__device__ double pow(double, double);
+__device__ float remainderf(float, float);
+__device__ double remainder(double, double);
+__device__ float remquof(float, float, int *);
+__device__ double remquo(double, double, int *);
+__device__ float rintf(float);
+__device__ double rint(double);
+__device__ float roundf(float);
+__device__ double round(double);
+__device__ float roundevenf(float);
+__device__ double roundeven(double);
+__device__ float scalblnf(float, long);
+__device__ double scalbln(double, long);
+__device__ float scalbnf(float, int);
+__device__ double scalbn(double, int);
+__device__ float sinf(float);
+__device__ double sin(double);
+__device__ void sincosf(float, float *, float *);
+__device__ void sincos(double, double *, double *);
+__device__ void sincospif(float, float *, float *);
+__device__ void sincospi(double, double *, double *);
+__device__ float sinhf(float);
+__device__ double sinh(double);
+__device__ float sqrtf(float);
+__device__ double sqrt(double);
+__device__ float tanf(float);
+__device__ double tan(double);
+__device__ float tanhf(float);
+__device__ double tanh(double);
+__device__ float tgammaf(float);
+__device__ double tgamma(double);
+__device__ float truncf(float);
+__device__ double trunc(double);
+}
+
+#endif // defined(__CUDA__) || defined(__HIP__)
+
+#endif // LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_MATH_H
diff --git a/offload/languages/kernel/CMakeLists.txt b/offload/languages/kernel/CMakeLists.txt
index a23e5727e6b8b..af33f94e1540c 100644
--- a/offload/languages/kernel/CMakeLists.txt
+++ b/offload/languages/kernel/CMakeLists.txt
@@ -107,6 +107,7 @@ install(FILES ${LLVM_OFFLOAD_KERNEL_PER_THREAD_DEFAULT_STREAM_OBJECT}
install(FILES
${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/DefineLanguageNames.inc
${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/LanguageErrors.h
+ ${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/LanguageMath.h
${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/LanguageRuntime.h
${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/UndefineLanguageNames.inc
${CMAKE_CURRENT_SOURCE_DIR}/include/LanguageLaunch.h
diff --git a/offload/test/offloading/language/math.cpp b/offload/test/offloading/language/math.cpp
new file mode 100644
index 0000000000000..b3e28ceb7346a
--- /dev/null
+++ b/offload/test/offloading/language/math.cpp
@@ -0,0 +1,107 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda
+// RUN: %t.cuda | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda -DOFFLOAD_TEST_LANGUAGE=cuda %s -o %t.cuda.omp -fopenmp
+// RUN: %t.cuda.omp | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip
+// RUN: %t.hip | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip -DOFFLOAD_TEST_LANGUAGE=hip %s -o %t.hip.omp -fopenmp
+// RUN: %t.hip.omp | %fcheck-generic
+// clang-format on
+
+// REQUIRES: gpu, libc
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+// clang-format off
+#include <stdio.h>
+#include "Inputs/DefineTestLanguageNames.inc"
+// clang-format on
+
+__global__ void math_kernel(double *D, float *F) {
+ D[0] = sqrt(D[0]);
+ D[1] = atan(D[1]);
+ D[2] = cos(D[2]);
+ D[3] = sin(D[3]);
+ D[4] = fabs(D[4]);
+ D[5] = floor(D[5]);
+ D[6] = ceil(D[6]);
+ D[7] = pow(D[7], D[8]);
+ D[8] = fma(D[8], 2.0, 1.0);
+ D[9] = copysign(D[9], D[0]);
+ D[10] = fmod(D[10], D[11]);
+ D[11] = exp2(D[12]);
+ D[12] = log2(D[11]);
+ D[13] = trunc(D[13]);
+ D[14] = round(D[14]);
+ D[15] = tan(D[15]);
+
+ F[0] = sqrtf(F[0]);
+ F[1] = atanf(F[1]);
+ F[2] = cosf(F[2]);
+ F[3] = sinf(F[3]);
+ F[4] = fabsf(F[4]);
+ F[5] = floorf(F[5]);
+ F[6] = ceilf(F[6]);
+ F[7] = powf(F[7], F[8]);
+ F[8] = fmaf(F[8], 2.0f, 1.0f);
+ F[9] = copysignf(F[9], F[0]);
+ F[10] = fmodf(F[10], F[11]);
+ F[11] = exp2f(F[12]);
+ F[12] = log2f(F[11]);
+ F[13] = truncf(F[13]);
+ F[14] = roundf(F[14]);
+ F[15] = tanf(F[15]);
+}
+
+int main(int argc, char **argv) {
+ double DHost[16] = {9.0, 0.0, 0.0, 0.0, -5.0, 1.25, 1.75, 2.0,
+ 3.0, -2.0, 5.0, 2.0, 1.0, 2.7, 2.4, 0.0};
+ float FHost[16] = {16.0f, 0.0f, 0.0f, 0.0f, -5.0f, 1.25f, 1.75f, 2.0f,
+ 3.0f, -2.0f, 5.0f, 2.0f, 1.0f, 2.7f, 2.4f, 0.0f};
+
+ double *D = nullptr;
+ float *F = nullptr;
+
+ if (Malloc(&D, sizeof(DHost)) != Success)
+ return 1;
+ if (Malloc(&F, sizeof(FHost)) != Success)
+ return 1;
+
+ if (Memcpy(D, DHost, sizeof(DHost), MemcpyHostToDevice) != Success)
+ return 1;
+ if (Memcpy(F, FHost, sizeof(FHost), MemcpyHostToDevice) != Success)
+ return 1;
+
+ math_kernel<<<1, 1>>>(D, F);
+ if (DeviceSynchronize() != Success)
+ return 1;
+
+ if (Memcpy(DHost, D, sizeof(DHost), MemcpyDeviceToHost) != Success)
+ return 1;
+ if (Memcpy(FHost, F, sizeof(FHost), MemcpyDeviceToHost) != Success)
+ return 1;
+
+ printf("double math A: %.1f %.1f %.1f %.1f %.1f %.1f %.1f %.1f\n", DHost[0],
+ DHost[1], DHost[2], DHost[3], DHost[4], DHost[5], DHost[6], DHost[7]);
+ // CHECK: double math A: 3.0 0.0 1.0 0.0 5.0 1.0 2.0 8.0
+ printf("double math B: %.1f %.1f %.1f %.1f %.1f %.1f %.1f %.1f\n", DHost[8],
+ DHost[9], DHost[10], DHost[11], DHost[12], DHost[13], DHost[14],
+ DHost[15]);
+ // CHECK: double math B: 7.0 2.0 1.0 2.0 1.0 2.0 2.0 0.0
+ printf("float math A: %.1f %.1f %.1f %.1f %.1f %.1f %.1f %.1f\n", FHost[0],
+ FHost[1], FHost[2], FHost[3], FHost[4], FHost[5], FHost[6], FHost[7]);
+ // CHECK: float math A: 4.0 0.0 1.0 0.0 5.0 1.0 2.0 8.0
+ printf("float math B: %.1f %.1f %.1f %.1f %.1f %.1f %.1f %.1f\n", FHost[8],
+ FHost[9], FHost[10], FHost[11], FHost[12], FHost[13], FHost[14],
+ FHost[15]);
+ // CHECK: float math B: 7.0 2.0 1.0 2.0 1.0 2.0 2.0 0.0
+
+ Free(D);
+ Free(F);
+}
More information about the llvm-branch-commits
mailing list