[libclc] [libclc] Move several 'native' builtins to CLC library (PR #129679)
Fraser Cormack via cfe-commits
cfe-commits at lists.llvm.org
Mon Mar 31 09:40:13 PDT 2025
https://github.com/frasercrmck updated https://github.com/llvm/llvm-project/pull/129679
>From aabdf579b3d88e59f8604943573c9beeafb4cf66 Mon Sep 17 00:00:00 2001
From: Fraser Cormack <fraser at codeplay.com>
Date: Tue, 4 Mar 2025 10:03:00 +0000
Subject: [PATCH] [libclc] Move several 'native' builtins to CLC library
This commit moves the 'native' builtins that use asm statements to
generate LLVM intrinsics to the CLC library. In doing so it converts
them to use the appropriate elementwise builtin to generate the same
intrinsic; there are no codegen changes to any target.
This work forms part of #127196 and indeed with this commit there are no
'generic' builtins using/abusing asm statements - the remaining builtins
are specific to the amdgpu and r600 targets.
---
libclc/amdgpu/lib/SOURCES | 4 --
libclc/clc/include/clc/math/clc_native_cos.h | 22 ++++++++++
libclc/clc/include/clc/math/clc_native_exp.h | 22 ++++++++++
libclc/clc/include/clc/math/clc_native_exp2.h | 22 ++++++++++
libclc/clc/include/clc/math/clc_native_log.h | 22 ++++++++++
.../clc/include/clc/math/clc_native_log10.h | 22 ++++++++++
libclc/clc/include/clc/math/clc_native_log2.h | 22 ++++++++++
libclc/clc/include/clc/math/clc_native_sin.h | 22 ++++++++++
libclc/clc/include/clc/math/clc_native_sqrt.h | 22 ++++++++++
libclc/clc/include/clc/math/unary_intrin.inc | 42 -------------------
libclc/clc/lib/amdgpu/SOURCES | 4 ++
.../lib/amdgpu/math/clc_native_exp.cl} | 6 ++-
.../lib/amdgpu/math/clc_native_exp.inc} | 4 +-
.../lib/amdgpu/math/clc_native_exp2.cl} | 6 +--
.../lib/amdgpu/math/clc_native_log.cl} | 6 ++-
.../lib/amdgpu/math/clc_native_log.inc} | 4 +-
.../lib/amdgpu/math/clc_native_log10.cl} | 6 ++-
.../lib/amdgpu/math/clc_native_log10.inc} | 4 +-
libclc/clc/lib/generic/SOURCES | 8 ++++
libclc/clc/lib/generic/math/clc_native_cos.cl | 16 +++++++
libclc/clc/lib/generic/math/clc_native_exp.cl | 16 +++++++
.../clc/lib/generic/math/clc_native_exp2.cl | 16 +++++++
libclc/clc/lib/generic/math/clc_native_log.cl | 16 +++++++
.../clc/lib/generic/math/clc_native_log10.cl | 16 +++++++
.../clc/lib/generic/math/clc_native_log2.cl | 16 +++++++
libclc/clc/lib/generic/math/clc_native_sin.cl | 16 +++++++
.../clc/lib/generic/math/clc_native_sqrt.cl | 16 +++++++
libclc/generic/lib/math/native_cos.cl | 7 ++--
libclc/generic/lib/math/native_exp.cl | 7 ++--
libclc/generic/lib/math/native_exp2.cl | 7 ++--
libclc/generic/lib/math/native_log.cl | 7 ++--
libclc/generic/lib/math/native_log10.cl | 7 ++--
libclc/generic/lib/math/native_log2.cl | 6 ++-
libclc/generic/lib/math/native_sin.cl | 7 ++--
libclc/generic/lib/math/native_sqrt.cl | 7 ++--
.../lib/math/native_unary_intrinsic.inc | 26 ------------
36 files changed, 369 insertions(+), 110 deletions(-)
create mode 100644 libclc/clc/include/clc/math/clc_native_cos.h
create mode 100644 libclc/clc/include/clc/math/clc_native_exp.h
create mode 100644 libclc/clc/include/clc/math/clc_native_exp2.h
create mode 100644 libclc/clc/include/clc/math/clc_native_log.h
create mode 100644 libclc/clc/include/clc/math/clc_native_log10.h
create mode 100644 libclc/clc/include/clc/math/clc_native_log2.h
create mode 100644 libclc/clc/include/clc/math/clc_native_sin.h
create mode 100644 libclc/clc/include/clc/math/clc_native_sqrt.h
delete mode 100644 libclc/clc/include/clc/math/unary_intrin.inc
rename libclc/{amdgpu/lib/math/native_exp.cl => clc/lib/amdgpu/math/clc_native_exp.cl} (74%)
rename libclc/{amdgpu/lib/math/native_exp.inc => clc/lib/amdgpu/math/clc_native_exp.inc} (75%)
rename libclc/{amdgpu/lib/math/native_exp2.cl => clc/lib/amdgpu/math/clc_native_exp2.cl} (72%)
rename libclc/{amdgpu/lib/math/native_log10.cl => clc/lib/amdgpu/math/clc_native_log.cl} (74%)
rename libclc/{amdgpu/lib/math/native_log10.inc => clc/lib/amdgpu/math/clc_native_log.inc} (74%)
rename libclc/{amdgpu/lib/math/native_log.cl => clc/lib/amdgpu/math/clc_native_log10.cl} (74%)
rename libclc/{amdgpu/lib/math/native_log.inc => clc/lib/amdgpu/math/clc_native_log10.inc} (73%)
create mode 100644 libclc/clc/lib/generic/math/clc_native_cos.cl
create mode 100644 libclc/clc/lib/generic/math/clc_native_exp.cl
create mode 100644 libclc/clc/lib/generic/math/clc_native_exp2.cl
create mode 100644 libclc/clc/lib/generic/math/clc_native_log.cl
create mode 100644 libclc/clc/lib/generic/math/clc_native_log10.cl
create mode 100644 libclc/clc/lib/generic/math/clc_native_log2.cl
create mode 100644 libclc/clc/lib/generic/math/clc_native_sin.cl
create mode 100644 libclc/clc/lib/generic/math/clc_native_sqrt.cl
delete mode 100644 libclc/generic/lib/math/native_unary_intrinsic.inc
diff --git a/libclc/amdgpu/lib/SOURCES b/libclc/amdgpu/lib/SOURCES
index ed5e45a37c18d..ab5da40711aa4 100644
--- a/libclc/amdgpu/lib/SOURCES
+++ b/libclc/amdgpu/lib/SOURCES
@@ -1,7 +1,3 @@
-math/native_exp.cl
-math/native_exp2.cl
-math/native_log.cl
-math/native_log10.cl
math/half_exp.cl
math/half_exp10.cl
math/half_exp2.cl
diff --git a/libclc/clc/include/clc/math/clc_native_cos.h b/libclc/clc/include/clc/math/clc_native_cos.h
new file mode 100644
index 0000000000000..8a580f13ad2aa
--- /dev/null
+++ b/libclc/clc/include/clc/math/clc_native_cos.h
@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache 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 __CLC_MATH_CLC_NATIVE_COS_H__
+#define __CLC_MATH_CLC_NATIVE_COS_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_cos
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_COS_H__
diff --git a/libclc/clc/include/clc/math/clc_native_exp.h b/libclc/clc/include/clc/math/clc_native_exp.h
new file mode 100644
index 0000000000000..48a1be616ea3e
--- /dev/null
+++ b/libclc/clc/include/clc/math/clc_native_exp.h
@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache 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 __CLC_MATH_CLC_NATIVE_EXP_H__
+#define __CLC_MATH_CLC_NATIVE_EXP_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_exp
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_EXP_H__
diff --git a/libclc/clc/include/clc/math/clc_native_exp2.h b/libclc/clc/include/clc/math/clc_native_exp2.h
new file mode 100644
index 0000000000000..bc0b32d6212b5
--- /dev/null
+++ b/libclc/clc/include/clc/math/clc_native_exp2.h
@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache 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 __CLC_MATH_CLC_NATIVE_EXP2_H__
+#define __CLC_MATH_CLC_NATIVE_EXP2_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_exp2
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_EXP2_H__
diff --git a/libclc/clc/include/clc/math/clc_native_log.h b/libclc/clc/include/clc/math/clc_native_log.h
new file mode 100644
index 0000000000000..ea0362503f670
--- /dev/null
+++ b/libclc/clc/include/clc/math/clc_native_log.h
@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache 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 __CLC_MATH_CLC_NATIVE_LOG_H__
+#define __CLC_MATH_CLC_NATIVE_LOG_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_log
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_LOG_H__
diff --git a/libclc/clc/include/clc/math/clc_native_log10.h b/libclc/clc/include/clc/math/clc_native_log10.h
new file mode 100644
index 0000000000000..c5cceeeba2952
--- /dev/null
+++ b/libclc/clc/include/clc/math/clc_native_log10.h
@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache 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 __CLC_MATH_CLC_NATIVE_LOG10_H__
+#define __CLC_MATH_CLC_NATIVE_LOG10_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_log10
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_LOG10_H__
diff --git a/libclc/clc/include/clc/math/clc_native_log2.h b/libclc/clc/include/clc/math/clc_native_log2.h
new file mode 100644
index 0000000000000..25375970cedc0
--- /dev/null
+++ b/libclc/clc/include/clc/math/clc_native_log2.h
@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache 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 __CLC_MATH_CLC_NATIVE_LOG2_H__
+#define __CLC_MATH_CLC_NATIVE_LOG2_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_log2
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_LOG2_H__
diff --git a/libclc/clc/include/clc/math/clc_native_sin.h b/libclc/clc/include/clc/math/clc_native_sin.h
new file mode 100644
index 0000000000000..878e19882ae6b
--- /dev/null
+++ b/libclc/clc/include/clc/math/clc_native_sin.h
@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache 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 __CLC_MATH_CLC_NATIVE_SIN_H__
+#define __CLC_MATH_CLC_NATIVE_SIN_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_sin
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_SIN_H__
diff --git a/libclc/clc/include/clc/math/clc_native_sqrt.h b/libclc/clc/include/clc/math/clc_native_sqrt.h
new file mode 100644
index 0000000000000..afff77c89b22f
--- /dev/null
+++ b/libclc/clc/include/clc/math/clc_native_sqrt.h
@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache 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 __CLC_MATH_CLC_NATIVE_SQRT_H__
+#define __CLC_MATH_CLC_NATIVE_SQRT_H__
+
+#define __FLOAT_ONLY
+#define __CLC_FUNCTION __clc_native_sqrt
+#define __CLC_BODY <clc/shared/unary_decl.inc>
+
+#include <clc/math/gentype.inc>
+
+#undef __CLC_BODY
+#undef __CLC_FUNCTION
+#undef __FLOAT_ONLY
+
+#endif // __CLC_MATH_CLC_NATIVE_SQRT_H__
diff --git a/libclc/clc/include/clc/math/unary_intrin.inc b/libclc/clc/include/clc/math/unary_intrin.inc
deleted file mode 100644
index 8028470114b8e..0000000000000
--- a/libclc/clc/include/clc/math/unary_intrin.inc
+++ /dev/null
@@ -1,42 +0,0 @@
-//===----------------------------------------------------------------------===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-#include <clc/clcfunc.h>
-#include <clc/clctypes.h>
-
-_CLC_OVERLOAD float __CLC_FUNCTION(float f) __asm(__CLC_INTRINSIC ".f32");
-_CLC_OVERLOAD float2 __CLC_FUNCTION(float2 f) __asm(__CLC_INTRINSIC ".v2f32");
-_CLC_OVERLOAD float3 __CLC_FUNCTION(float3 f) __asm(__CLC_INTRINSIC ".v3f32");
-_CLC_OVERLOAD float4 __CLC_FUNCTION(float4 f) __asm(__CLC_INTRINSIC ".v4f32");
-_CLC_OVERLOAD float8 __CLC_FUNCTION(float8 f) __asm(__CLC_INTRINSIC ".v8f32");
-_CLC_OVERLOAD float16 __CLC_FUNCTION(float16 f) __asm(__CLC_INTRINSIC
- ".v16f32");
-
-#ifdef cl_khr_fp64
-#pragma OPENCL EXTENSION cl_khr_fp64 : enable
-_CLC_OVERLOAD double __CLC_FUNCTION(double d) __asm(__CLC_INTRINSIC ".f64");
-_CLC_OVERLOAD double2 __CLC_FUNCTION(double2 d) __asm(__CLC_INTRINSIC ".v2f64");
-_CLC_OVERLOAD double3 __CLC_FUNCTION(double3 d) __asm(__CLC_INTRINSIC ".v3f64");
-_CLC_OVERLOAD double4 __CLC_FUNCTION(double4 d) __asm(__CLC_INTRINSIC ".v4f64");
-_CLC_OVERLOAD double8 __CLC_FUNCTION(double8 d) __asm(__CLC_INTRINSIC ".v8f64");
-_CLC_OVERLOAD double16 __CLC_FUNCTION(double16 d) __asm(__CLC_INTRINSIC
- ".v16f64");
-#endif
-
-#ifdef cl_khr_fp16
-#pragma OPENCL EXTENSION cl_khr_fp16 : enable
-_CLC_OVERLOAD half __CLC_FUNCTION(half d) __asm(__CLC_INTRINSIC ".f16");
-_CLC_OVERLOAD half2 __CLC_FUNCTION(half2 d) __asm(__CLC_INTRINSIC ".v2f16");
-_CLC_OVERLOAD half3 __CLC_FUNCTION(half3 d) __asm(__CLC_INTRINSIC ".v3f16");
-_CLC_OVERLOAD half4 __CLC_FUNCTION(half4 d) __asm(__CLC_INTRINSIC ".v4f16");
-_CLC_OVERLOAD half8 __CLC_FUNCTION(half8 d) __asm(__CLC_INTRINSIC ".v8f16");
-_CLC_OVERLOAD half16 __CLC_FUNCTION(half16 d) __asm(__CLC_INTRINSIC ".v16f16");
-#endif
-
-#undef __CLC_FUNCTION
-#undef __CLC_INTRINSIC
diff --git a/libclc/clc/lib/amdgpu/SOURCES b/libclc/clc/lib/amdgpu/SOURCES
index fd64a862021e8..607603c1e6fa7 100644
--- a/libclc/clc/lib/amdgpu/SOURCES
+++ b/libclc/clc/lib/amdgpu/SOURCES
@@ -1 +1,5 @@
+math/clc_native_exp2.cl
+math/clc_native_exp.cl
+math/clc_native_log10.cl
+math/clc_native_log.cl
math/clc_sqrt_fp64.cl
diff --git a/libclc/amdgpu/lib/math/native_exp.cl b/libclc/clc/lib/amdgpu/math/clc_native_exp.cl
similarity index 74%
rename from libclc/amdgpu/lib/math/native_exp.cl
rename to libclc/clc/lib/amdgpu/math/clc_native_exp.cl
index e62b79d4ec9fa..591ecb0ac00b5 100644
--- a/libclc/amdgpu/lib/math/native_exp.cl
+++ b/libclc/clc/lib/amdgpu/math/clc_native_exp.cl
@@ -6,8 +6,10 @@
//
//===----------------------------------------------------------------------===//
-#include <clc/clc.h>
+#include <clc/float/definitions.h>
+#include <clc/internal/clc.h>
+#include <clc/math/clc_native_exp2.h>
-#define __CLC_BODY <native_exp.inc>
+#define __CLC_BODY <clc_native_exp.inc>
#define __FLOAT_ONLY
#include <clc/math/gentype.inc>
diff --git a/libclc/amdgpu/lib/math/native_exp.inc b/libclc/clc/lib/amdgpu/math/clc_native_exp.inc
similarity index 75%
rename from libclc/amdgpu/lib/math/native_exp.inc
rename to libclc/clc/lib/amdgpu/math/clc_native_exp.inc
index d7dbd888d2988..cf2d48c4054ed 100644
--- a/libclc/amdgpu/lib/math/native_exp.inc
+++ b/libclc/clc/lib/amdgpu/math/clc_native_exp.inc
@@ -6,6 +6,6 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE native_exp(__CLC_GENTYPE val) {
- return native_exp2(val * M_LOG2E_F);
+_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_native_exp(__CLC_GENTYPE val) {
+ return __clc_native_exp2(val * M_LOG2E_F);
}
diff --git a/libclc/amdgpu/lib/math/native_exp2.cl b/libclc/clc/lib/amdgpu/math/clc_native_exp2.cl
similarity index 72%
rename from libclc/amdgpu/lib/math/native_exp2.cl
rename to libclc/clc/lib/amdgpu/math/clc_native_exp2.cl
index 39ae914b19634..76b1850fce574 100644
--- a/libclc/amdgpu/lib/math/native_exp2.cl
+++ b/libclc/clc/lib/amdgpu/math/clc_native_exp2.cl
@@ -6,11 +6,11 @@
//
//===----------------------------------------------------------------------===//
-#include <clc/clc.h>
#include <clc/clcmacro.h>
+#include <clc/internal/clc.h>
-_CLC_OVERLOAD _CLC_DEF float native_exp2(float val) {
+_CLC_OVERLOAD _CLC_DEF float __clc_native_exp2(float val) {
return __builtin_amdgcn_exp2f(val);
}
-_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, native_exp2, float)
+_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_native_exp2, float)
diff --git a/libclc/amdgpu/lib/math/native_log10.cl b/libclc/clc/lib/amdgpu/math/clc_native_log.cl
similarity index 74%
rename from libclc/amdgpu/lib/math/native_log10.cl
rename to libclc/clc/lib/amdgpu/math/clc_native_log.cl
index 7cbe1f98988d5..fdff8cd65f344 100644
--- a/libclc/amdgpu/lib/math/native_log10.cl
+++ b/libclc/clc/lib/amdgpu/math/clc_native_log.cl
@@ -6,8 +6,10 @@
//
//===----------------------------------------------------------------------===//
-#include <clc/clc.h>
+#include <clc/float/definitions.h>
+#include <clc/internal/clc.h>
+#include <clc/math/clc_native_log2.h>
-#define __CLC_BODY <native_log10.inc>
+#define __CLC_BODY <clc_native_log.inc>
#define __FLOAT_ONLY
#include <clc/math/gentype.inc>
diff --git a/libclc/amdgpu/lib/math/native_log10.inc b/libclc/clc/lib/amdgpu/math/clc_native_log.inc
similarity index 74%
rename from libclc/amdgpu/lib/math/native_log10.inc
rename to libclc/clc/lib/amdgpu/math/clc_native_log.inc
index 2aef43a8ed1a5..91d76d249e94c 100644
--- a/libclc/amdgpu/lib/math/native_log10.inc
+++ b/libclc/clc/lib/amdgpu/math/clc_native_log.inc
@@ -6,6 +6,6 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE native_log10(__CLC_GENTYPE val) {
- return native_log2(val) * (M_LN2_F / M_LN10_F);
+_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_native_log(__CLC_GENTYPE val) {
+ return __clc_native_log2(val) * (1.0f / M_LOG2E_F);
}
diff --git a/libclc/amdgpu/lib/math/native_log.cl b/libclc/clc/lib/amdgpu/math/clc_native_log10.cl
similarity index 74%
rename from libclc/amdgpu/lib/math/native_log.cl
rename to libclc/clc/lib/amdgpu/math/clc_native_log10.cl
index ce3258bb5b37b..0668a635d24d7 100644
--- a/libclc/amdgpu/lib/math/native_log.cl
+++ b/libclc/clc/lib/amdgpu/math/clc_native_log10.cl
@@ -6,8 +6,10 @@
//
//===----------------------------------------------------------------------===//
-#include <clc/clc.h>
+#include <clc/float/definitions.h>
+#include <clc/internal/clc.h>
+#include <clc/math/clc_native_log2.h>
-#define __CLC_BODY <native_log.inc>
+#define __CLC_BODY <clc_native_log10.inc>
#define __FLOAT_ONLY
#include <clc/math/gentype.inc>
diff --git a/libclc/amdgpu/lib/math/native_log.inc b/libclc/clc/lib/amdgpu/math/clc_native_log10.inc
similarity index 73%
rename from libclc/amdgpu/lib/math/native_log.inc
rename to libclc/clc/lib/amdgpu/math/clc_native_log10.inc
index 820e4929f02cf..c91d698609793 100644
--- a/libclc/amdgpu/lib/math/native_log.inc
+++ b/libclc/clc/lib/amdgpu/math/clc_native_log10.inc
@@ -6,6 +6,6 @@
//
//===----------------------------------------------------------------------===//
-_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE native_log(__CLC_GENTYPE val) {
- return native_log2(val) * (1.0f / M_LOG2E_F);
+_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_native_log10(__CLC_GENTYPE val) {
+ return __clc_native_log2(val) * (M_LN2_F / M_LN10_F);
}
diff --git a/libclc/clc/lib/generic/SOURCES b/libclc/clc/lib/generic/SOURCES
index 3c648edd52a7c..a7c488f94ccc8 100644
--- a/libclc/clc/lib/generic/SOURCES
+++ b/libclc/clc/lib/generic/SOURCES
@@ -45,6 +45,14 @@ math/clc_log2.cl
math/clc_mad.cl
math/clc_modf.cl
math/clc_nan.cl
+math/clc_native_cos.cl
+math/clc_native_exp.cl
+math/clc_native_exp2.cl
+math/clc_native_log.cl
+math/clc_native_log10.cl
+math/clc_native_log2.cl
+math/clc_native_sin.cl
+math/clc_native_sqrt.cl
math/clc_nextafter.cl
math/clc_pow.cl
math/clc_pown.cl
diff --git a/libclc/clc/lib/generic/math/clc_native_cos.cl b/libclc/clc/lib/generic/math/clc_native_cos.cl
new file mode 100644
index 0000000000000..de56fdec48d24
--- /dev/null
+++ b/libclc/clc/lib/generic/math/clc_native_cos.cl
@@ -0,0 +1,16 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <clc/internal/clc.h>
+
+#define __FLOAT_ONLY
+#define FUNCTION __clc_native_cos
+#define __CLC_FUNCTION(x) __builtin_elementwise_cos
+#define __CLC_BODY <clc/shared/unary_def.inc>
+
+#include <clc/math/gentype.inc>
diff --git a/libclc/clc/lib/generic/math/clc_native_exp.cl b/libclc/clc/lib/generic/math/clc_native_exp.cl
new file mode 100644
index 0000000000000..400270a6163a4
--- /dev/null
+++ b/libclc/clc/lib/generic/math/clc_native_exp.cl
@@ -0,0 +1,16 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <clc/internal/clc.h>
+
+#define __FLOAT_ONLY
+#define FUNCTION __clc_native_exp
+#define __CLC_FUNCTION(x) __builtin_elementwise_exp
+#define __CLC_BODY <clc/shared/unary_def.inc>
+
+#include <clc/math/gentype.inc>
diff --git a/libclc/clc/lib/generic/math/clc_native_exp2.cl b/libclc/clc/lib/generic/math/clc_native_exp2.cl
new file mode 100644
index 0000000000000..427d901fcdb19
--- /dev/null
+++ b/libclc/clc/lib/generic/math/clc_native_exp2.cl
@@ -0,0 +1,16 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <clc/internal/clc.h>
+
+#define __FLOAT_ONLY
+#define FUNCTION __clc_native_exp2
+#define __CLC_FUNCTION(x) __builtin_elementwise_exp2
+#define __CLC_BODY <clc/shared/unary_def.inc>
+
+#include <clc/math/gentype.inc>
diff --git a/libclc/clc/lib/generic/math/clc_native_log.cl b/libclc/clc/lib/generic/math/clc_native_log.cl
new file mode 100644
index 0000000000000..85f188b654282
--- /dev/null
+++ b/libclc/clc/lib/generic/math/clc_native_log.cl
@@ -0,0 +1,16 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <clc/internal/clc.h>
+
+#define __FLOAT_ONLY
+#define FUNCTION __clc_native_log
+#define __CLC_FUNCTION(x) __builtin_elementwise_log
+#define __CLC_BODY <clc/shared/unary_def.inc>
+
+#include <clc/math/gentype.inc>
diff --git a/libclc/clc/lib/generic/math/clc_native_log10.cl b/libclc/clc/lib/generic/math/clc_native_log10.cl
new file mode 100644
index 0000000000000..624018e4481bf
--- /dev/null
+++ b/libclc/clc/lib/generic/math/clc_native_log10.cl
@@ -0,0 +1,16 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <clc/internal/clc.h>
+
+#define __FLOAT_ONLY
+#define FUNCTION __clc_native_log10
+#define __CLC_FUNCTION(x) __builtin_elementwise_log10
+#define __CLC_BODY <clc/shared/unary_def.inc>
+
+#include <clc/math/gentype.inc>
diff --git a/libclc/clc/lib/generic/math/clc_native_log2.cl b/libclc/clc/lib/generic/math/clc_native_log2.cl
new file mode 100644
index 0000000000000..2c8c18e61ca5d
--- /dev/null
+++ b/libclc/clc/lib/generic/math/clc_native_log2.cl
@@ -0,0 +1,16 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <clc/internal/clc.h>
+
+#define __FLOAT_ONLY
+#define FUNCTION __clc_native_log2
+#define __CLC_FUNCTION(x) __builtin_elementwise_log2
+#define __CLC_BODY <clc/shared/unary_def.inc>
+
+#include <clc/math/gentype.inc>
diff --git a/libclc/clc/lib/generic/math/clc_native_sin.cl b/libclc/clc/lib/generic/math/clc_native_sin.cl
new file mode 100644
index 0000000000000..22b988bf4375f
--- /dev/null
+++ b/libclc/clc/lib/generic/math/clc_native_sin.cl
@@ -0,0 +1,16 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <clc/internal/clc.h>
+
+#define __FLOAT_ONLY
+#define FUNCTION __clc_native_sin
+#define __CLC_FUNCTION(x) __builtin_elementwise_sin
+#define __CLC_BODY <clc/shared/unary_def.inc>
+
+#include <clc/math/gentype.inc>
diff --git a/libclc/clc/lib/generic/math/clc_native_sqrt.cl b/libclc/clc/lib/generic/math/clc_native_sqrt.cl
new file mode 100644
index 0000000000000..ed022ef1fee1f
--- /dev/null
+++ b/libclc/clc/lib/generic/math/clc_native_sqrt.cl
@@ -0,0 +1,16 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <clc/internal/clc.h>
+
+#define __FLOAT_ONLY
+#define FUNCTION __clc_native_sqrt
+#define __CLC_FUNCTION(x) __builtin_elementwise_sqrt
+#define __CLC_BODY <clc/shared/unary_def.inc>
+
+#include <clc/math/gentype.inc>
diff --git a/libclc/generic/lib/math/native_cos.cl b/libclc/generic/lib/math/native_cos.cl
index d81a067f11044..0f4d46c1d73e4 100644
--- a/libclc/generic/lib/math/native_cos.cl
+++ b/libclc/generic/lib/math/native_cos.cl
@@ -7,9 +7,10 @@
//===----------------------------------------------------------------------===//
#include <clc/clc.h>
+#include <clc/math/clc_native_cos.h>
-#define __CLC_NATIVE_INTRINSIC cos
-
-#define __CLC_BODY <native_unary_intrinsic.inc>
#define __FLOAT_ONLY
+#define FUNCTION native_cos
+#define __CLC_BODY <clc/shared/unary_def.inc>
+
#include <clc/math/gentype.inc>
diff --git a/libclc/generic/lib/math/native_exp.cl b/libclc/generic/lib/math/native_exp.cl
index 8f4531343415c..c1d08ec2c7e1f 100644
--- a/libclc/generic/lib/math/native_exp.cl
+++ b/libclc/generic/lib/math/native_exp.cl
@@ -7,9 +7,10 @@
//===----------------------------------------------------------------------===//
#include <clc/clc.h>
+#include <clc/math/clc_native_exp.h>
-#define __CLC_NATIVE_INTRINSIC exp
-
-#define __CLC_BODY <native_unary_intrinsic.inc>
#define __FLOAT_ONLY
+#define FUNCTION native_exp
+#define __CLC_BODY <clc/shared/unary_def.inc>
+
#include <clc/math/gentype.inc>
diff --git a/libclc/generic/lib/math/native_exp2.cl b/libclc/generic/lib/math/native_exp2.cl
index ecde4a6761d22..ceb570733b974 100644
--- a/libclc/generic/lib/math/native_exp2.cl
+++ b/libclc/generic/lib/math/native_exp2.cl
@@ -7,9 +7,10 @@
//===----------------------------------------------------------------------===//
#include <clc/clc.h>
+#include <clc/math/clc_native_exp2.h>
-#define __CLC_NATIVE_INTRINSIC exp2
-
-#define __CLC_BODY <native_unary_intrinsic.inc>
#define __FLOAT_ONLY
+#define FUNCTION native_exp2
+#define __CLC_BODY <clc/shared/unary_def.inc>
+
#include <clc/math/gentype.inc>
diff --git a/libclc/generic/lib/math/native_log.cl b/libclc/generic/lib/math/native_log.cl
index 5731e09d21c9c..adc2ff495f8b5 100644
--- a/libclc/generic/lib/math/native_log.cl
+++ b/libclc/generic/lib/math/native_log.cl
@@ -7,9 +7,10 @@
//===----------------------------------------------------------------------===//
#include <clc/clc.h>
+#include <clc/math/clc_native_log.h>
-#define __CLC_NATIVE_INTRINSIC log
-
-#define __CLC_BODY <native_unary_intrinsic.inc>
#define __FLOAT_ONLY
+#define FUNCTION native_log
+#define __CLC_BODY <clc/shared/unary_def.inc>
+
#include <clc/math/gentype.inc>
diff --git a/libclc/generic/lib/math/native_log10.cl b/libclc/generic/lib/math/native_log10.cl
index eab7a6f14d035..f63292124f3b7 100644
--- a/libclc/generic/lib/math/native_log10.cl
+++ b/libclc/generic/lib/math/native_log10.cl
@@ -7,9 +7,10 @@
//===----------------------------------------------------------------------===//
#include <clc/clc.h>
+#include <clc/math/clc_native_log10.h>
-#define __CLC_NATIVE_INTRINSIC log10
-
-#define __CLC_BODY <native_unary_intrinsic.inc>
#define __FLOAT_ONLY
+#define FUNCTION native_log10
+#define __CLC_BODY <clc/shared/unary_def.inc>
+
#include <clc/math/gentype.inc>
diff --git a/libclc/generic/lib/math/native_log2.cl b/libclc/generic/lib/math/native_log2.cl
index 0db4be0b5e083..6b079872b1e0a 100644
--- a/libclc/generic/lib/math/native_log2.cl
+++ b/libclc/generic/lib/math/native_log2.cl
@@ -7,8 +7,10 @@
//===----------------------------------------------------------------------===//
#include <clc/clc.h>
+#include <clc/math/clc_native_log2.h>
-#define __CLC_NATIVE_INTRINSIC log2
-#define __CLC_BODY <native_unary_intrinsic.inc>
#define __FLOAT_ONLY
+#define FUNCTION native_log2
+#define __CLC_BODY <clc/shared/unary_def.inc>
+
#include <clc/math/gentype.inc>
diff --git a/libclc/generic/lib/math/native_sin.cl b/libclc/generic/lib/math/native_sin.cl
index 0e2ced09fa2dd..50265b3936272 100644
--- a/libclc/generic/lib/math/native_sin.cl
+++ b/libclc/generic/lib/math/native_sin.cl
@@ -7,9 +7,10 @@
//===----------------------------------------------------------------------===//
#include <clc/clc.h>
+#include <clc/math/clc_native_sin.h>
-#define __CLC_NATIVE_INTRINSIC sin
-
-#define __CLC_BODY <native_unary_intrinsic.inc>
#define __FLOAT_ONLY
+#define FUNCTION native_sin
+#define __CLC_BODY <clc/shared/unary_def.inc>
+
#include <clc/math/gentype.inc>
diff --git a/libclc/generic/lib/math/native_sqrt.cl b/libclc/generic/lib/math/native_sqrt.cl
index 1b668e5976ef7..4cd022e8bbeba 100644
--- a/libclc/generic/lib/math/native_sqrt.cl
+++ b/libclc/generic/lib/math/native_sqrt.cl
@@ -7,9 +7,10 @@
//===----------------------------------------------------------------------===//
#include <clc/clc.h>
+#include <clc/math/clc_native_sqrt.h>
-#define __CLC_NATIVE_INTRINSIC sqrt
-
-#define __CLC_BODY <native_unary_intrinsic.inc>
#define __FLOAT_ONLY
+#define FUNCTION native_sqrt
+#define __CLC_BODY <clc/shared/unary_def.inc>
+
#include <clc/math/gentype.inc>
diff --git a/libclc/generic/lib/math/native_unary_intrinsic.inc b/libclc/generic/lib/math/native_unary_intrinsic.inc
deleted file mode 100644
index c118ec095692f..0000000000000
--- a/libclc/generic/lib/math/native_unary_intrinsic.inc
+++ /dev/null
@@ -1,26 +0,0 @@
-//===----------------------------------------------------------------------===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-#include <clc/utils.h>
-
-#ifdef __CLC_SCALAR
-#define __CLC_FUNCTION __CLC_XCONCAT(__clc_native_, __CLC_NATIVE_INTRINSIC)
-#define __CLC_INTRINSIC "llvm." __CLC_XSTR(__CLC_NATIVE_INTRINSIC)
-
-#undef cl_khr_fp64
-#include <clc/math/unary_intrin.inc>
-
-#endif
-
-#define __CLC_FUNCTION __CLC_XCONCAT(native_, __CLC_NATIVE_INTRINSIC)
-
-_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __CLC_FUNCTION(__CLC_GENTYPE val) {
- return __CLC_XCONCAT(__clc_native_, __CLC_NATIVE_INTRINSIC)(val);
-}
-
-#undef __CLC_FUNCTION
More information about the cfe-commits
mailing list