[libclc] Move __clc_ldexp to CLC library

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 → 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 → 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/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

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

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

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