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 +#include + +#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 #include +#include +#include #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 - -#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 #include #include +#include #include #include #include #include #include -_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_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 #include -#include -#include +#include -_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