[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.

Author: frasercrmck

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;
 )

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

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__

libclc/clc/include/clc/math/clc_ldexp.inc

@@ -0,0 +1 @@
+_CLC_DECL _CLC_OVERLOAD __CLC_GENTYPE __clc_ldexp(__CLC_GENTYPE, __CLC_INTN);

libclc/clc/lib/amdgcn/SOURCES

@@ -0,0 +1 @@
+math/clc_ldexp_override.cl

libclc/amdgcn/lib/math/ldexp.cl renamed to 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

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

libclc/generic/lib/math/clc_ldexp.cl renamed to 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

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

libclc/generic/include/math/clc_ldexp.h

@@ -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

libclc/generic/lib/SOURCES

@@ -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

libclc/generic/lib/math/ldexp.cl

@@ -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

libclc/generic/lib/math/ldexp.inc

@@ -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