[libclc] [libclc] Move __clc_ldexp to CLC library (PR #126078)
Fraser Cormack via cfe-commits
cfe-commits at lists.llvm.org
Thu Feb 6 06:54:12 PST 2025
https://github.com/frasercrmck created https://github.com/llvm/llvm-project/pull/126078
This function was already conceptually in the CLC namespace - this just formally moves it over.
Note however that this commit marks a change in how libclc functions may be overridden by targets.
Until now we have been using a purely build-system-based approach where targets could register identically-named files which took responsibility for the implementation of the builtin in its entirety.
This system wasn't well equipped to deal with AMD's overriding of __clc_ldexp for only a subset of types, and furthermore conditionally on a pre-defined macro.
One option for handling this would be to require AMD to duplicate code for the versions of __clc_ldexp it's *not* interested in overriding. We could also make it easier for targets to re-define CLC functions through macros or .inc files. Both of these have obvious downsides. We could also keep AMD's overriding in the OpenCL layer and bypass CLC altogether, but this has limited use.
We could use weak linkage on the "base" implementations of CLC functions, and allow targets to opt-in to providing their own implementations on a much finer granularity. This commit supports this as a proof of concept; we could expand it to all CLC builtins if accepted.
Note that the existing filename-based "claiming" approach is still in effect, so targets have to name their overrides differently to have both files compiled. This could also be refined.
>From dc4bcfc33b6e1668f6480af881df0529d2bb7725 Mon Sep 17 00:00:00 2001
From: Fraser Cormack <fraser at codeplay.com>
Date: Wed, 5 Feb 2025 10:38:46 +0000
Subject: [PATCH] [libclc] Move __clc_ldexp to CLC library
This function was already conceptually in the CLC namespace - this just
formally moves it over.
Note however that this commit marks a change in how libclc functions may
be overridden by targets.
Until now we have been using a purely build-system-based approach where
targets could register identically-named files which took responsibility
for the implementation of the builtin in its entirety.
This system wasn't well equipped to deal with AMD's overriding of
__clc_ldexp for only a subset of types, and furthermore conditionally on
a pre-defined macro.
One option for handling this would be to require AMD to duplicate code
for the versions of __clc_ldexp it's *not* interested in overriding. We
could also make it easier for targets to re-define CLC functions through
macros or .inc files. Both of these have obvious downsides. We could
also keep AMD's overriding in the OpenCL layer and bypass CLC
altogether, but this has limited use.
We could use weak linkage on the "base" implementations of CLC
functions, and allow targets to opt-in to providing their own
implementations on a much finer granulaity. This commit supports this as
a proof of concept; we could expand it to all CLC builtins if accepted.
Note that the existing filename-based "claiming" approach is still in
effect, so targets have to name their overrides differently to have both
files compiled. This could also be refined.
---
libclc/CMakeLists.txt | 1 +
libclc/amdgcn/lib/SOURCES | 1 -
libclc/clc/include/clc/math/clc_ldexp.h | 7 +++
libclc/clc/include/clc/math/clc_ldexp.inc | 1 +
libclc/clc/include/clc/relational/clc_isinf.h | 7 ---
libclc/clc/lib/amdgcn/SOURCES | 1 +
.../lib/amdgcn/math/clc_ldexp_override.cl} | 25 ++++-------
libclc/clc/lib/clspv/SOURCES | 2 +
libclc/clc/lib/generic/SOURCES | 1 +
.../lib => clc/lib/generic}/math/clc_ldexp.cl | 43 +++++++++++--------
libclc/clc/lib/spirv/SOURCES | 3 ++
libclc/clspv/lib/SOURCES | 1 -
libclc/generic/include/math/clc_ldexp.h | 13 ------
libclc/generic/lib/SOURCES | 1 -
libclc/generic/lib/math/ldexp.cl | 12 +++---
libclc/generic/lib/math/ldexp.inc | 7 +--
libclc/spirv/lib/SOURCES | 1 -
17 files changed, 56 insertions(+), 71 deletions(-)
create mode 100644 libclc/clc/include/clc/math/clc_ldexp.h
create mode 100644 libclc/clc/include/clc/math/clc_ldexp.inc
create mode 100644 libclc/clc/lib/amdgcn/SOURCES
rename libclc/{amdgcn/lib/math/ldexp.cl => clc/lib/amdgcn/math/clc_ldexp_override.cl} (70%)
rename libclc/{generic/lib => clc/lib/generic}/math/clc_ldexp.cl (76%)
delete mode 100644 libclc/generic/include/math/clc_ldexp.h
diff --git a/libclc/CMakeLists.txt b/libclc/CMakeLists.txt
index ff52153354e0a9c..9e65c6340296ed7 100644
--- a/libclc/CMakeLists.txt
+++ b/libclc/CMakeLists.txt
@@ -32,6 +32,7 @@ set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS
spirv64/lib/SOURCES;
# CLC internal libraries
clc/lib/generic/SOURCES;
+ clc/lib/amdgcn/SOURCES;
clc/lib/clspv/SOURCES;
clc/lib/clspv64/SOURCES;
clc/lib/spirv/SOURCES;
diff --git a/libclc/amdgcn/lib/SOURCES b/libclc/amdgcn/lib/SOURCES
index 4ea66385fe50ee3..6c6e77db0d84b56 100644
--- a/libclc/amdgcn/lib/SOURCES
+++ b/libclc/amdgcn/lib/SOURCES
@@ -1,7 +1,6 @@
cl_khr_int64_extended_atomics/minmax_helpers.ll
math/fmax.cl
math/fmin.cl
-math/ldexp.cl
mem_fence/fence.cl
synchronization/barrier.cl
workitem/get_global_offset.cl
diff --git a/libclc/clc/include/clc/math/clc_ldexp.h b/libclc/clc/include/clc/math/clc_ldexp.h
new file mode 100644
index 000000000000000..8a639fa03f04851
--- /dev/null
+++ b/libclc/clc/include/clc/math/clc_ldexp.h
@@ -0,0 +1,7 @@
+#ifndef __CLC_MATH_CLC_LDEXP_H__
+#define __CLC_MATH_CLC_LDEXP_H__
+
+#define __CLC_BODY <clc/math/clc_ldexp.inc>
+#include <clc/math/gentype.inc>
+
+#endif // __CLC_MATH_CLC_LDEXP_H__
diff --git a/libclc/clc/include/clc/math/clc_ldexp.inc b/libclc/clc/include/clc/math/clc_ldexp.inc
new file mode 100644
index 000000000000000..759b3270d7544d8
--- /dev/null
+++ b/libclc/clc/include/clc/math/clc_ldexp.inc
@@ -0,0 +1 @@
+_CLC_DECL _CLC_OVERLOAD __CLC_GENTYPE __clc_ldexp(__CLC_GENTYPE, __CLC_INTN);
diff --git a/libclc/clc/include/clc/relational/clc_isinf.h b/libclc/clc/include/clc/relational/clc_isinf.h
index 3f60bec5654a2e2..b666953d4a8e6c7 100644
--- a/libclc/clc/include/clc/relational/clc_isinf.h
+++ b/libclc/clc/include/clc/relational/clc_isinf.h
@@ -1,11 +1,6 @@
#ifndef __CLC_RELATIONAL_CLC_ISINF_H__
#define __CLC_RELATIONAL_CLC_ISINF_H__
-#if defined(CLC_CLSPV) || defined(CLC_SPIRV)
-// clspv and spir-v targets provide their own OpenCL-compatible isinf
-#define __clc_isinf isinf
-#else
-
#include <clc/clcfunc.h>
#include <clc/clctypes.h>
@@ -37,6 +32,4 @@ _CLC_VECTOR_ISINF_DECL(short, half)
#undef _CLC_ISINF_DECL
#undef _CLC_VECTOR_ISINF_DECL
-#endif
-
#endif // __CLC_RELATIONAL_CLC_ISINF_H__
diff --git a/libclc/clc/lib/amdgcn/SOURCES b/libclc/clc/lib/amdgcn/SOURCES
new file mode 100644
index 000000000000000..b85f80d3b29bbcf
--- /dev/null
+++ b/libclc/clc/lib/amdgcn/SOURCES
@@ -0,0 +1 @@
+math/clc_ldexp_override.cl
diff --git a/libclc/amdgcn/lib/math/ldexp.cl b/libclc/clc/lib/amdgcn/math/clc_ldexp_override.cl
similarity index 70%
rename from libclc/amdgcn/lib/math/ldexp.cl
rename to libclc/clc/lib/amdgcn/math/clc_ldexp_override.cl
index d46d2dc606818a6..03324db2bf48d43 100644
--- a/libclc/amdgcn/lib/math/ldexp.cl
+++ b/libclc/clc/lib/amdgcn/math/clc_ldexp_override.cl
@@ -20,27 +20,18 @@
* THE SOFTWARE.
*/
-#include <clc/clc.h>
#include <clc/clcmacro.h>
+#include <clc/internal/clc.h>
+#include <clc/math/clc_ldexp.h>
#ifdef __HAS_LDEXPF__
-#define BUILTINF __builtin_amdgcn_ldexpf
-#else
-#include "math/clc_ldexp.h"
-#define BUILTINF __clc_ldexp
-#endif
-
// This defines all the ldexp(floatN, intN) variants.
-_CLC_DEFINE_BINARY_BUILTIN(float, ldexp, BUILTINF, float, int);
+_CLC_DEFINE_BINARY_BUILTIN(float, __clc_ldexp, __builtin_amdgcn_ldexpf, float, int);
+#endif
#ifdef cl_khr_fp64
- #pragma OPENCL EXTENSION cl_khr_fp64 : enable
- // This defines all the ldexp(doubleN, intN) variants.
- _CLC_DEFINE_BINARY_BUILTIN(double, ldexp, __builtin_amdgcn_ldexp, double, int);
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+// This defines all the ldexp(doubleN, intN) variants.
+_CLC_DEFINE_BINARY_BUILTIN(double, __clc_ldexp, __builtin_amdgcn_ldexp, double,
+ int);
#endif
-
-// This defines all the ldexp(GENTYPE, int);
-#define __CLC_BODY <../../../generic/lib/math/ldexp.inc>
-#include <clc/math/gentype.inc>
-
-#undef BUILTINF
diff --git a/libclc/clc/lib/clspv/SOURCES b/libclc/clc/lib/clspv/SOURCES
index 2fe07f62a328ca0..7cbe34302fa435a 100644
--- a/libclc/clc/lib/clspv/SOURCES
+++ b/libclc/clc/lib/clspv/SOURCES
@@ -14,10 +14,12 @@
../generic/math/clc_copysign.cl
../generic/math/clc_fabs.cl
../generic/math/clc_floor.cl
+../generic/math/clc_ldexp.cl
../generic/math/clc_mad.cl
../generic/math/clc_nextafter.cl
../generic/math/clc_rint.cl
../generic/math/clc_trunc.cl
+../generic/relational/clc_isinf.cl
../generic/relational/clc_isnan.cl
../generic/relational/clc_select.cl
../generic/shared/clc_clamp.cl
diff --git a/libclc/clc/lib/generic/SOURCES b/libclc/clc/lib/generic/SOURCES
index 2d31f9ac23fec27..e58408366009462 100644
--- a/libclc/clc/lib/generic/SOURCES
+++ b/libclc/clc/lib/generic/SOURCES
@@ -20,6 +20,7 @@ math/clc_ceil.cl
math/clc_copysign.cl
math/clc_fabs.cl
math/clc_floor.cl
+math/clc_ldexp.cl
math/clc_mad.cl
math/clc_nextafter.cl
math/clc_rint.cl
diff --git a/libclc/generic/lib/math/clc_ldexp.cl b/libclc/clc/lib/generic/math/clc_ldexp.cl
similarity index 76%
rename from libclc/generic/lib/math/clc_ldexp.cl
rename to libclc/clc/lib/generic/math/clc_ldexp.cl
index a03f05d21775611..ad54755bee2768f 100644
--- a/libclc/generic/lib/math/clc_ldexp.cl
+++ b/libclc/clc/lib/generic/math/clc_ldexp.cl
@@ -20,21 +20,22 @@
* THE SOFTWARE.
*/
-#include <clc/clc.h>
#include <clc/clcmacro.h>
#include <clc/integer/clc_add_sat.h>
+#include <clc/internal/clc.h>
#include <clc/math/clc_subnormal_config.h>
#include <clc/math/math.h>
#include <clc/relational/clc_isinf.h>
#include <clc/relational/clc_isnan.h>
#include <clc/shared/clc_clamp.h>
-_CLC_DEF _CLC_OVERLOAD float __clc_ldexp(float x, int n) {
+#define _CLC_DEF_ldexp _CLC_DEF __attribute__((weak))
- if (!__clc_fp32_subnormals_supported()) {
+_CLC_DEF_ldexp _CLC_OVERLOAD float __clc_ldexp(float x, int n) {
+ if (!__clc_fp32_subnormals_supported()) {
// This treats subnormals as zeros
- int i = as_int(x);
+ int i = __clc_as_int(x);
int e = (i >> 23) & 0xff;
int m = i & 0x007fffff;
int s = i & 0x80000000;
@@ -45,7 +46,7 @@ _CLC_DEF _CLC_OVERLOAD float __clc_ldexp(float x, int n) {
mr = c ? m : mr;
int er = c ? e : v;
er = e ? er : e;
- return as_float(s | (er << 23) | mr);
+ return __clc_as_float(s | (er << 23) | mr);
}
/* supports denormal values */
@@ -54,7 +55,7 @@ _CLC_DEF _CLC_OVERLOAD float __clc_ldexp(float x, int n) {
uint val_ui;
uint sign;
int exponent;
- val_ui = as_uint(x);
+ val_ui = __clc_as_uint(x);
sign = val_ui & 0x80000000;
val_ui = val_ui & 0x7fffffff; /* remove the sign bit */
int val_x = val_ui;
@@ -64,7 +65,9 @@ _CLC_DEF _CLC_OVERLOAD float __clc_ldexp(float x, int n) {
/* denormal support */
int fbh =
- 127 - (as_uint((float)(as_float(val_ui | 0x3f800000) - 1.0f)) >> 23);
+ 127 -
+ (__clc_as_uint((float)(__clc_as_float(val_ui | 0x3f800000) - 1.0f)) >>
+ 23);
int dexponent = 25 - fbh;
uint dval_ui = (((val_ui << fbh) & 0x007fffff) | (dexponent << 23));
int ex = dexponent + n - multiplier;
@@ -77,7 +80,7 @@ _CLC_DEF _CLC_OVERLOAD float __clc_ldexp(float x, int n) {
dval_ui = dexponent > 254 ? 0x7f800000 : dval_ui; /*overflow*/
dval_ui = dexponent < -multiplier ? 0 : dval_ui; /*underflow*/
dval_ui = dval_ui | sign;
- val_f = as_float(dval_ui);
+ val_f = __clc_as_float(dval_ui);
exponent += n;
@@ -91,22 +94,24 @@ _CLC_DEF _CLC_OVERLOAD float __clc_ldexp(float x, int n) {
val_ui = val_ui | sign;
val_ui = dexp == 0 ? dval_ui : val_ui;
- val_f = as_float(val_ui);
+ val_f = __clc_as_float(val_ui);
val_f = __clc_isnan(x) | __clc_isinf(x) | val_x == 0 ? x : val_f;
return val_f;
}
+_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF_ldexp, float, __clc_ldexp, float, int);
+
#ifdef cl_khr_fp64
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
-_CLC_DEF _CLC_OVERLOAD double __clc_ldexp(double x, int n) {
- long l = as_ulong(x);
+_CLC_DEF_ldexp _CLC_OVERLOAD double __clc_ldexp(double x, int n) {
+ long l = __clc_as_ulong(x);
int e = (l >> 52) & 0x7ff;
long s = l & 0x8000000000000000;
- ulong ux = as_ulong(x * 0x1.0p+53);
+ ulong ux = __clc_as_ulong(x * 0x1.0p+53);
int de = ((int)(ux >> 52) & 0x7ff) - 53;
int c = e == 0;
e = c ? de : e;
@@ -118,28 +123,30 @@ _CLC_DEF _CLC_OVERLOAD double __clc_ldexp(double x, int n) {
ux &= ~EXPBITS_DP64;
- double mr = as_double(ux | ((ulong)(v + 53) << 52));
+ double mr = __clc_as_double(ux | ((ulong)(v + 53) << 52));
mr = mr * 0x1.0p-53;
- mr = v > 0 ? as_double(ux | ((ulong)v << 52)) : mr;
+ mr = v > 0 ? __clc_as_double(ux | ((ulong)v << 52)) : mr;
- mr = v == 0x7ff ? as_double(s | PINFBITPATT_DP64) : mr;
- mr = v < -53 ? as_double(s) : mr;
+ mr = v == 0x7ff ? __clc_as_double(s | PINFBITPATT_DP64) : mr;
+ mr = v < -53 ? __clc_as_double(s) : mr;
mr = ((n == 0) | __clc_isinf(x) | (x == 0)) ? x : mr;
return mr;
}
+_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF_ldexp, double, __clc_ldexp, double, int);
+
#endif
#ifdef cl_khr_fp16
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
-_CLC_OVERLOAD _CLC_DEF half __clc_ldexp(half x, int n) {
+_CLC_OVERLOAD _CLC_DEF_ldexp half __clc_ldexp(half x, int n) {
return (half)__clc_ldexp((float)x, n);
}
-_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, half, __clc_ldexp, half, int);
+_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF_ldexp, half, __clc_ldexp, half, int);
#endif
diff --git a/libclc/clc/lib/spirv/SOURCES b/libclc/clc/lib/spirv/SOURCES
index 96040a3aebd8355..2a53ff7c74c4c07 100644
--- a/libclc/clc/lib/spirv/SOURCES
+++ b/libclc/clc/lib/spirv/SOURCES
@@ -18,9 +18,12 @@
../generic/math/clc_copysign.cl
../generic/math/clc_fabs.cl
../generic/math/clc_floor.cl
+../generic/math/clc_ldexp.cl
../generic/math/clc_mad.cl
../generic/math/clc_nextafter.cl
../generic/math/clc_rint.cl
../generic/math/clc_trunc.cl
+../generic/relational/clc_isinf.cl
+../generic/relational/clc_isnan.cl
../generic/relational/clc_select.cl
../generic/shared/clc_clamp.cl
diff --git a/libclc/clspv/lib/SOURCES b/libclc/clspv/lib/SOURCES
index 0d6091ce20e4495..15c437beb03b9a6 100644
--- a/libclc/clspv/lib/SOURCES
+++ b/libclc/clspv/lib/SOURCES
@@ -19,7 +19,6 @@ subnormal_config.cl
../../generic/lib/math/clc_exp10.cl
../../generic/lib/math/clc_fmod.cl
../../generic/lib/math/clc_hypot.cl
-../../generic/lib/math/clc_ldexp.cl
../../generic/lib/math/clc_pow.cl
../../generic/lib/math/clc_pown.cl
../../generic/lib/math/clc_powr.cl
diff --git a/libclc/generic/include/math/clc_ldexp.h b/libclc/generic/include/math/clc_ldexp.h
deleted file mode 100644
index eb83f16240185e1..000000000000000
--- a/libclc/generic/include/math/clc_ldexp.h
+++ /dev/null
@@ -1,13 +0,0 @@
-#include <clc/clcfunc.h>
-
-_CLC_DEF _CLC_OVERLOAD float __clc_ldexp(float, int);
-
-#ifdef cl_khr_fp64
-#pragma OPENCL EXTENSION cl_khr_fp64 : enable
-_CLC_DEF _CLC_OVERLOAD double __clc_ldexp(double, int);
-#endif
-
-#ifdef cl_khr_fp16
-#pragma OPENCL EXTENSION cl_khr_fp16 : enable
-_CLC_DEF _CLC_OVERLOAD half __clc_ldexp(half, int);
-#endif
diff --git a/libclc/generic/lib/SOURCES b/libclc/generic/lib/SOURCES
index a62c87902a6a7e7..7aba3d6693ea359 100644
--- a/libclc/generic/lib/SOURCES
+++ b/libclc/generic/lib/SOURCES
@@ -131,7 +131,6 @@ math/half_tan.cl
math/clc_hypot.cl
math/hypot.cl
math/ilogb.cl
-math/clc_ldexp.cl
math/ldexp.cl
math/lgamma.cl
math/lgamma_r.cl
diff --git a/libclc/generic/lib/math/ldexp.cl b/libclc/generic/lib/math/ldexp.cl
index 72708f74b5cc1c1..56c8d1243bbf51e 100644
--- a/libclc/generic/lib/math/ldexp.cl
+++ b/libclc/generic/lib/math/ldexp.cl
@@ -20,26 +20,26 @@
* THE SOFTWARE.
*/
-#include "math/clc_ldexp.h"
#include <clc/clc.h>
#include <clc/clcmacro.h>
-#include <clc/math/clc_subnormal_config.h>
-#include <clc/math/math.h>
+#include <clc/math/clc_ldexp.h>
-_CLC_DEFINE_BINARY_BUILTIN(float, ldexp, __clc_ldexp, float, int)
+_CLC_DEFINE_BINARY_BUILTIN_NO_SCALARIZE(float, ldexp, __clc_ldexp, float, int)
#ifdef cl_khr_fp64
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
-_CLC_DEFINE_BINARY_BUILTIN(double, ldexp, __clc_ldexp, double, int)
+_CLC_DEFINE_BINARY_BUILTIN_NO_SCALARIZE(double, ldexp, __clc_ldexp, double, int)
+
#endif
#ifdef cl_khr_fp16
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
-_CLC_DEFINE_BINARY_BUILTIN(half, ldexp, __clc_ldexp, half, int)
+_CLC_DEFINE_BINARY_BUILTIN_NO_SCALARIZE(half, ldexp, __clc_ldexp, half, int)
+
#endif
// This defines all the ldexp(GENTYPE, int) variants
diff --git a/libclc/generic/lib/math/ldexp.inc b/libclc/generic/lib/math/ldexp.inc
index d6144d7cb6d6207..1547eb7cc0252b8 100644
--- a/libclc/generic/lib/math/ldexp.inc
+++ b/libclc/generic/lib/math/ldexp.inc
@@ -20,15 +20,10 @@
* THE SOFTWARE.
*/
-// TODO: Enable half precision when ldexp is implemented.
-#if __CLC_FPSIZE > 16
-
#ifndef __CLC_SCALAR
_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE ldexp(__CLC_GENTYPE x, int n) {
- return ldexp(x, (__CLC_INTN)n);
+ return __clc_ldexp(x, (__CLC_INTN)n);
}
#endif
-
-#endif
diff --git a/libclc/spirv/lib/SOURCES b/libclc/spirv/lib/SOURCES
index 854cba614c8bfd8..06462195674489c 100644
--- a/libclc/spirv/lib/SOURCES
+++ b/libclc/spirv/lib/SOURCES
@@ -52,7 +52,6 @@ math/fma.cl
../../generic/lib/math/clc_hypot.cl
../../generic/lib/math/hypot.cl
../../generic/lib/math/ilogb.cl
-../../generic/lib/math/clc_ldexp.cl
../../generic/lib/math/ldexp.cl
../../generic/lib/math/lgamma.cl
../../generic/lib/math/lgamma_r.cl
More information about the cfe-commits
mailing list