[libclc] d5038b3 - [libclc] Move __clc_ldexp to CLC library (#126078)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Feb 26 03:20:29 PST 2025
Author: Fraser Cormack
Date: 2025-02-26T11:20:25Z
New Revision: d5038b3774485d617e1300cf2f7b98c2460b9042
URL: https://github.com/llvm/llvm-project/commit/d5038b3774485d617e1300cf2f7b98c2460b9042
DIFF: https://github.com/llvm/llvm-project/commit/d5038b3774485d617e1300cf2f7b98c2460b9042.diff
LOG: [libclc] Move __clc_ldexp to CLC library (#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.
Added:
libclc/clc/include/clc/math/clc_ldexp.h
libclc/clc/include/clc/math/clc_ldexp.inc
libclc/clc/lib/amdgcn/SOURCES
libclc/clc/lib/amdgcn/math/clc_ldexp_override.cl
libclc/clc/lib/generic/math/clc_ldexp.cl
Modified:
libclc/CMakeLists.txt
libclc/amdgcn/lib/SOURCES
libclc/clc/lib/generic/SOURCES
libclc/clspv/lib/SOURCES
libclc/generic/lib/SOURCES
libclc/generic/lib/math/ldexp.cl
libclc/generic/lib/math/ldexp.inc
libclc/spirv/lib/SOURCES
Removed:
libclc/amdgcn/lib/math/ldexp.cl
libclc/generic/include/math/clc_ldexp.h
libclc/generic/lib/math/clc_ldexp.cl
################################################################################
diff --git a/libclc/CMakeLists.txt b/libclc/CMakeLists.txt
index 22c3075801785..c1a1dd42bcb46 100644
--- a/libclc/CMakeLists.txt
+++ b/libclc/CMakeLists.txt
@@ -28,6 +28,7 @@ set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS
spirv/lib/SOURCES;
# CLC internal libraries
clc/lib/generic/SOURCES;
+ clc/lib/amdgcn/SOURCES;
clc/lib/clspv/SOURCES;
clc/lib/spirv/SOURCES;
)
diff --git a/libclc/amdgcn/lib/SOURCES b/libclc/amdgcn/lib/SOURCES
index 4ea66385fe50e..6c6e77db0d84b 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 0000000000000..8a639fa03f048
--- /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 0000000000000..759b3270d7544
--- /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/lib/amdgcn/SOURCES b/libclc/clc/lib/amdgcn/SOURCES
new file mode 100644
index 0000000000000..b85f80d3b29bb
--- /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 d46d2dc606818..03324db2bf48d 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/generic/SOURCES b/libclc/clc/lib/generic/SOURCES
index dbf3c1bd31a0d..1d16cd5e2d18a 100644
--- a/libclc/clc/lib/generic/SOURCES
+++ b/libclc/clc/lib/generic/SOURCES
@@ -23,6 +23,7 @@ math/clc_fabs.cl
math/clc_fma.cl
math/clc_floor.cl
math/clc_frexp.cl
+math/clc_ldexp.cl
math/clc_log.cl
math/clc_log10.cl
math/clc_log2.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 a03f05d217756..ad54755bee276 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/clspv/lib/SOURCES b/libclc/clspv/lib/SOURCES
index 0d6091ce20e44..15c437beb03b9 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 eb83f16240185..0000000000000
--- 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 bea78adf2cf08..eb14fd84c96b3 100644
--- a/libclc/generic/lib/SOURCES
+++ b/libclc/generic/lib/SOURCES
@@ -130,7 +130,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 72708f74b5cc1..56c8d1243bbf5 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 d6144d7cb6d62..1547eb7cc0252 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 82416a9e69bfc..35eef0f364707 100644
--- a/libclc/spirv/lib/SOURCES
+++ b/libclc/spirv/lib/SOURCES
@@ -51,7 +51,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