https://github.com/wenju-he created https://github.com/llvm/llvm-project/pull/145678
With this PR, if we have customized implementation for scalar or vector length = 2, we don't need to write new macros, e.g. https://github.com/intel/llvm/blob/fb18321705f6/libclc/clc/include/clc/clcmacro.h#L15 Undef __HALF_ONLY, __FLOAT_ONLY and __DOUBLE_ONLY at the end of clc/include/clc/math/gentype.inc llvm-diff shows no change to nvptx64--nvidiacl.bc and amdgcn--amdhsa.bc >From 4e6aa2a55e4519fb3c21b2351766a3cfc6772b68 Mon Sep 17 00:00:00 2001 From: Wenju He <wenju...@intel.com> Date: Wed, 25 Jun 2025 13:24:56 +0200 Subject: [PATCH] [NFC][libclc] Refactor _CLC_*_VECTORIZE macros to functions in .inc files With this PR, if we have customized implementation for scalar or vector length = 2, we don't need to write new macros, e.g. https://github.com/intel/llvm/blob/fb18321705f6/libclc/clc/include/clc/clcmacro.h#L15 Undef __HALF_ONLY, __FLOAT_ONLY and __DOUBLE_ONLY at the end of clc/include/clc/math/gentype.inc llvm-diff shows no change to nvptx64--nvidiacl.bc and amdgcn--amdhsa.bc --- libclc/clc/include/clc/clcmacro.h | 113 -------------- .../include/clc/geometric/clc_fast_distance.h | 1 - .../include/clc/geometric/clc_fast_length.h | 1 - .../clc/geometric/clc_fast_normalize.h | 1 - libclc/clc/include/clc/math/clc_exp_helper.h | 2 - libclc/clc/include/clc/math/clc_half_cos.h | 1 - libclc/clc/include/clc/math/clc_half_divide.h | 1 - libclc/clc/include/clc/math/clc_half_exp.h | 1 - libclc/clc/include/clc/math/clc_half_exp10.h | 1 - libclc/clc/include/clc/math/clc_half_exp2.h | 1 - libclc/clc/include/clc/math/clc_half_log.h | 1 - libclc/clc/include/clc/math/clc_half_log10.h | 1 - libclc/clc/include/clc/math/clc_half_log2.h | 1 - libclc/clc/include/clc/math/clc_half_powr.h | 1 - libclc/clc/include/clc/math/clc_half_recip.h | 1 - libclc/clc/include/clc/math/clc_half_rsqrt.h | 1 - libclc/clc/include/clc/math/clc_half_sin.h | 1 - libclc/clc/include/clc/math/clc_half_sqrt.h | 1 - libclc/clc/include/clc/math/clc_half_tan.h | 1 - libclc/clc/include/clc/math/clc_native_cos.h | 1 - .../clc/include/clc/math/clc_native_divide.h | 1 - libclc/clc/include/clc/math/clc_native_exp.h | 1 - .../clc/include/clc/math/clc_native_exp10.h | 1 - libclc/clc/include/clc/math/clc_native_exp2.h | 1 - libclc/clc/include/clc/math/clc_native_log.h | 1 - .../clc/include/clc/math/clc_native_log10.h | 1 - libclc/clc/include/clc/math/clc_native_log2.h | 1 - libclc/clc/include/clc/math/clc_native_powr.h | 1 - .../clc/include/clc/math/clc_native_recip.h | 1 - .../clc/include/clc/math/clc_native_rsqrt.h | 1 - libclc/clc/include/clc/math/clc_native_sin.h | 1 - libclc/clc/include/clc/math/clc_native_sqrt.h | 1 - libclc/clc/include/clc/math/clc_native_tan.h | 1 - .../clc/include/clc/math/clc_sincos_helpers.h | 4 - libclc/clc/include/clc/math/gentype.inc | 4 + .../clc/shared/binary_def_scalarize.inc | 121 +++++++++++++++ .../clc/shared/ternary_def_scalarize.inc | 143 ++++++++++++++++++ .../clc/shared/unary_def_scalarize.inc | 87 +++++++++++ libclc/clc/lib/amdgcn/math/clc_fmax.cl | 8 +- libclc/clc/lib/amdgcn/math/clc_fmin.cl | 8 +- .../clc/lib/amdgcn/math/clc_ldexp_override.cl | 17 ++- libclc/clc/lib/amdgpu/math/clc_native_exp2.cl | 11 +- libclc/clc/lib/amdgpu/math/clc_sqrt_fp64.cl | 5 +- libclc/clc/lib/clspv/math/clc_sw_fma.cl | 6 +- libclc/clc/lib/generic/integer/clc_clz.cl | 11 +- libclc/clc/lib/generic/integer/clc_ctz.cl | 11 +- libclc/clc/lib/generic/math/clc_erf.cl | 12 +- libclc/clc/lib/generic/math/clc_erfc.cl | 12 +- libclc/clc/lib/generic/math/clc_fmax.cl | 28 +++- libclc/clc/lib/generic/math/clc_fmin.cl | 27 +++- libclc/clc/lib/generic/math/clc_fmod.cl | 16 +- libclc/clc/lib/generic/math/clc_ldexp.cl | 12 +- libclc/clc/lib/generic/math/clc_log.cl | 10 +- libclc/clc/lib/generic/math/clc_log10.cl | 12 +- libclc/clc/lib/generic/math/clc_log2.cl | 12 +- libclc/clc/lib/generic/math/clc_remainder.cl | 17 ++- .../lib/generic/math/clc_sincos_helpers.cl | 4 - libclc/clc/lib/generic/math/clc_sw_fma.cl | 6 +- libclc/clc/lib/generic/math/clc_tgamma.cl | 12 +- libclc/clc/lib/r600/math/clc_fmax.cl | 15 +- libclc/clc/lib/r600/math/clc_fmin.cl | 15 +- libclc/clc/lib/r600/math/clc_native_rsqrt.cl | 5 +- .../clc/lib/r600/math/clc_rsqrt_override.cl | 12 +- .../clc/opencl/geometric/fast_distance.h | 1 - .../clc/opencl/geometric/fast_length.h | 1 - .../clc/opencl/geometric/fast_normalize.h | 1 - .../opencl/include/clc/opencl/math/half_cos.h | 1 - .../opencl/include/clc/opencl/math/half_exp.h | 1 - .../include/clc/opencl/math/half_exp10.h | 1 - .../include/clc/opencl/math/half_exp2.h | 1 - .../opencl/include/clc/opencl/math/half_log.h | 1 - .../include/clc/opencl/math/half_log10.h | 1 - .../include/clc/opencl/math/half_log2.h | 1 - .../include/clc/opencl/math/half_recip.h | 1 - .../include/clc/opencl/math/half_rsqrt.h | 1 - .../opencl/include/clc/opencl/math/half_sin.h | 1 - .../include/clc/opencl/math/half_sqrt.h | 1 - .../opencl/include/clc/opencl/math/half_tan.h | 1 - .../include/clc/opencl/math/native_cos.h | 1 - .../include/clc/opencl/math/native_exp.h | 1 - .../include/clc/opencl/math/native_exp10.h | 1 - .../include/clc/opencl/math/native_exp2.h | 1 - .../include/clc/opencl/math/native_log.h | 1 - .../include/clc/opencl/math/native_log10.h | 1 - .../include/clc/opencl/math/native_log2.h | 1 - .../include/clc/opencl/math/native_recip.h | 1 - .../include/clc/opencl/math/native_rsqrt.h | 1 - .../include/clc/opencl/math/native_sin.h | 1 - .../include/clc/opencl/math/native_sqrt.h | 1 - .../include/clc/opencl/math/native_tan.h | 1 - 90 files changed, 560 insertions(+), 276 deletions(-) create mode 100644 libclc/clc/include/clc/shared/binary_def_scalarize.inc create mode 100644 libclc/clc/include/clc/shared/ternary_def_scalarize.inc create mode 100644 libclc/clc/include/clc/shared/unary_def_scalarize.inc diff --git a/libclc/clc/include/clc/clcmacro.h b/libclc/clc/include/clc/clcmacro.h index b712fe5cf326c..5c67c937cb1d4 100644 --- a/libclc/clc/include/clc/clcmacro.h +++ b/libclc/clc/include/clc/clcmacro.h @@ -12,111 +12,6 @@ #include <clc/internal/clc.h> #include <clc/utils.h> -#define _CLC_UNARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE) \ - DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x) { \ - return (RET_TYPE##2)(FUNCTION(x.s0), FUNCTION(x.s1)); \ - } \ - \ - DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x) { \ - return (RET_TYPE##3)(FUNCTION(x.s0), FUNCTION(x.s1), FUNCTION(x.s2)); \ - } \ - \ - DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x) { \ - return (RET_TYPE##4)(FUNCTION(x.s0), FUNCTION(x.s1), FUNCTION(x.s2), \ - FUNCTION(x.s3)); \ - } \ - \ - DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x) { \ - return (RET_TYPE##8)(FUNCTION(x.s0), FUNCTION(x.s1), FUNCTION(x.s2), \ - FUNCTION(x.s3), FUNCTION(x.s4), FUNCTION(x.s5), \ - FUNCTION(x.s6), FUNCTION(x.s7)); \ - } \ - \ - DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x) { \ - return (RET_TYPE##16)( \ - FUNCTION(x.s0), FUNCTION(x.s1), FUNCTION(x.s2), FUNCTION(x.s3), \ - FUNCTION(x.s4), FUNCTION(x.s5), FUNCTION(x.s6), FUNCTION(x.s7), \ - FUNCTION(x.s8), FUNCTION(x.s9), FUNCTION(x.sa), FUNCTION(x.sb), \ - FUNCTION(x.sc), FUNCTION(x.sd), FUNCTION(x.se), FUNCTION(x.sf)); \ - } - -#define _CLC_BINARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, \ - ARG2_TYPE) \ - DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x, ARG2_TYPE##2 y) { \ - return (RET_TYPE##2)(FUNCTION(x.s0, y.s0), FUNCTION(x.s1, y.s1)); \ - } \ - \ - DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x, ARG2_TYPE##3 y) { \ - return (RET_TYPE##3)(FUNCTION(x.s0, y.s0), FUNCTION(x.s1, y.s1), \ - FUNCTION(x.s2, y.s2)); \ - } \ - \ - DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x, ARG2_TYPE##4 y) { \ - return (RET_TYPE##4)(FUNCTION(x.s0, y.s0), FUNCTION(x.s1, y.s1), \ - FUNCTION(x.s2, y.s2), FUNCTION(x.s3, y.s3)); \ - } \ - \ - DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x, ARG2_TYPE##8 y) { \ - return (RET_TYPE##8)(FUNCTION(x.s0, y.s0), FUNCTION(x.s1, y.s1), \ - FUNCTION(x.s2, y.s2), FUNCTION(x.s3, y.s3), \ - FUNCTION(x.s4, y.s4), FUNCTION(x.s5, y.s5), \ - FUNCTION(x.s6, y.s6), FUNCTION(x.s7, y.s7)); \ - } \ - \ - DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x, ARG2_TYPE##16 y) { \ - return (RET_TYPE##16)( \ - FUNCTION(x.s0, y.s0), FUNCTION(x.s1, y.s1), FUNCTION(x.s2, y.s2), \ - FUNCTION(x.s3, y.s3), FUNCTION(x.s4, y.s4), FUNCTION(x.s5, y.s5), \ - FUNCTION(x.s6, y.s6), FUNCTION(x.s7, y.s7), FUNCTION(x.s8, y.s8), \ - FUNCTION(x.s9, y.s9), FUNCTION(x.sa, y.sa), FUNCTION(x.sb, y.sb), \ - FUNCTION(x.sc, y.sc), FUNCTION(x.sd, y.sd), FUNCTION(x.se, y.se), \ - FUNCTION(x.sf, y.sf)); \ - } - -#define _CLC_TERNARY_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, \ - ARG2_TYPE, ARG3_TYPE) \ - DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE##2 x, ARG2_TYPE##2 y, \ - ARG3_TYPE##2 z) { \ - return (RET_TYPE##2)(FUNCTION(x.s0, y.s0, z.s0), \ - FUNCTION(x.s1, y.s1, z.s1)); \ - } \ - \ - DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x, ARG2_TYPE##3 y, \ - ARG3_TYPE##3 z) { \ - return (RET_TYPE##3)(FUNCTION(x.s0, y.s0, z.s0), \ - FUNCTION(x.s1, y.s1, z.s1), \ - FUNCTION(x.s2, y.s2, z.s2)); \ - } \ - \ - DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x, ARG2_TYPE##4 y, \ - ARG3_TYPE##4 z) { \ - return (RET_TYPE##4)( \ - FUNCTION(x.s0, y.s0, z.s0), FUNCTION(x.s1, y.s1, z.s1), \ - FUNCTION(x.s2, y.s2, z.s2), FUNCTION(x.s3, y.s3, z.s3)); \ - } \ - \ - DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x, ARG2_TYPE##8 y, \ - ARG3_TYPE##8 z) { \ - return (RET_TYPE##8)( \ - FUNCTION(x.s0, y.s0, z.s0), FUNCTION(x.s1, y.s1, z.s1), \ - FUNCTION(x.s2, y.s2, z.s2), FUNCTION(x.s3, y.s3, z.s3), \ - FUNCTION(x.s4, y.s4, z.s4), FUNCTION(x.s5, y.s5, z.s5), \ - FUNCTION(x.s6, y.s6, z.s6), FUNCTION(x.s7, y.s7, z.s7)); \ - } \ - \ - DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x, ARG2_TYPE##16 y, \ - ARG3_TYPE##16 z) { \ - return (RET_TYPE##16)( \ - FUNCTION(x.s0, y.s0, z.s0), FUNCTION(x.s1, y.s1, z.s1), \ - FUNCTION(x.s2, y.s2, z.s2), FUNCTION(x.s3, y.s3, z.s3), \ - FUNCTION(x.s4, y.s4, z.s4), FUNCTION(x.s5, y.s5, z.s5), \ - FUNCTION(x.s6, y.s6, z.s6), FUNCTION(x.s7, y.s7, z.s7), \ - FUNCTION(x.s8, y.s8, z.s8), FUNCTION(x.s9, y.s9, z.s9), \ - FUNCTION(x.sa, y.sa, z.sa), FUNCTION(x.sb, y.sb, z.sb), \ - FUNCTION(x.sc, y.sc, z.sc), FUNCTION(x.sd, y.sd, z.sd), \ - FUNCTION(x.se, y.se, z.se), FUNCTION(x.sf, y.sf, z.sf)); \ - } - #define _CLC_V_V_VP_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, \ ADDR_SPACE, ARG2_TYPE) \ DECLSPEC __CLC_XCONCAT(RET_TYPE, 2) \ @@ -171,12 +66,4 @@ FUNCTION(x.sf, ptr + 15)); \ } -#define _CLC_DEFINE_BINARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, \ - ARG2_TYPE) \ - _CLC_DEF _CLC_OVERLOAD RET_TYPE FUNCTION(ARG1_TYPE x, ARG2_TYPE y) { \ - return BUILTIN(x, y); \ - } \ - _CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE, \ - ARG2_TYPE) - #endif // __CLC_CLCMACRO_H__ diff --git a/libclc/clc/include/clc/geometric/clc_fast_distance.h b/libclc/clc/include/clc/geometric/clc_fast_distance.h index 489bea4299516..17165e56367ca 100644 --- a/libclc/clc/include/clc/geometric/clc_fast_distance.h +++ b/libclc/clc/include/clc/geometric/clc_fast_distance.h @@ -15,7 +15,6 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION #endif // __CLC_GEOMETRIC_CLC_FAST_DISTANCE_H__ diff --git a/libclc/clc/include/clc/geometric/clc_fast_length.h b/libclc/clc/include/clc/geometric/clc_fast_length.h index 40552634fff94..a2d533e79c0ef 100644 --- a/libclc/clc/include/clc/geometric/clc_fast_length.h +++ b/libclc/clc/include/clc/geometric/clc_fast_length.h @@ -15,7 +15,6 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION #endif // __CLC_GEOMETRIC_CLC_FAST_LENGTH_H__ diff --git a/libclc/clc/include/clc/geometric/clc_fast_normalize.h b/libclc/clc/include/clc/geometric/clc_fast_normalize.h index 66eed8b83ab18..a168ac1d46783 100644 --- a/libclc/clc/include/clc/geometric/clc_fast_normalize.h +++ b/libclc/clc/include/clc/geometric/clc_fast_normalize.h @@ -17,6 +17,5 @@ #undef __CLC_FUNCTION #undef __CLC_GEOMETRIC_RET_GENTYPE -#undef __FLOAT_ONLY #endif // __CLC_GEOMETRIC_CLC_FAST_NORMALIZE_H__ diff --git a/libclc/clc/include/clc/math/clc_exp_helper.h b/libclc/clc/include/clc/math/clc_exp_helper.h index 89c725ca0d5c9..0c5028bd23858 100644 --- a/libclc/clc/include/clc/math/clc_exp_helper.h +++ b/libclc/clc/include/clc/math/clc_exp_helper.h @@ -14,6 +14,4 @@ #include <clc/math/gentype.inc> -#undef __DOUBLE_ONLY - #endif // __CLC_MATH_CLC_EXP_HELPER diff --git a/libclc/clc/include/clc/math/clc_half_cos.h b/libclc/clc/include/clc/math/clc_half_cos.h index d0efc45793af9..4c70572c99070 100644 --- a/libclc/clc/include/clc/math/clc_half_cos.h +++ b/libclc/clc/include/clc/math/clc_half_cos.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_HALF_COS_H__ diff --git a/libclc/clc/include/clc/math/clc_half_divide.h b/libclc/clc/include/clc/math/clc_half_divide.h index 5d1e7b9e33e39..d87577e3bc201 100644 --- a/libclc/clc/include/clc/math/clc_half_divide.h +++ b/libclc/clc/include/clc/math/clc_half_divide.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_HALF_DIVIDE_H__ diff --git a/libclc/clc/include/clc/math/clc_half_exp.h b/libclc/clc/include/clc/math/clc_half_exp.h index e95f3d42085c8..a66f9c2ab16cb 100644 --- a/libclc/clc/include/clc/math/clc_half_exp.h +++ b/libclc/clc/include/clc/math/clc_half_exp.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_HALF_EXP_H__ diff --git a/libclc/clc/include/clc/math/clc_half_exp10.h b/libclc/clc/include/clc/math/clc_half_exp10.h index e4cce6e63bf7a..a61720af4d1b5 100644 --- a/libclc/clc/include/clc/math/clc_half_exp10.h +++ b/libclc/clc/include/clc/math/clc_half_exp10.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_HALF_EXP10_H__ diff --git a/libclc/clc/include/clc/math/clc_half_exp2.h b/libclc/clc/include/clc/math/clc_half_exp2.h index 6fc99735f0928..ca96d25f77c58 100644 --- a/libclc/clc/include/clc/math/clc_half_exp2.h +++ b/libclc/clc/include/clc/math/clc_half_exp2.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_HALF_EXP2_H__ diff --git a/libclc/clc/include/clc/math/clc_half_log.h b/libclc/clc/include/clc/math/clc_half_log.h index e44e686499663..95bf1a9745605 100644 --- a/libclc/clc/include/clc/math/clc_half_log.h +++ b/libclc/clc/include/clc/math/clc_half_log.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_HALF_LOG_H__ diff --git a/libclc/clc/include/clc/math/clc_half_log10.h b/libclc/clc/include/clc/math/clc_half_log10.h index 23e0ba116c728..359e37ce6cf18 100644 --- a/libclc/clc/include/clc/math/clc_half_log10.h +++ b/libclc/clc/include/clc/math/clc_half_log10.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_HALF_LOG10_H__ diff --git a/libclc/clc/include/clc/math/clc_half_log2.h b/libclc/clc/include/clc/math/clc_half_log2.h index 8ea2439eb3db6..82bfa9bd6afa6 100644 --- a/libclc/clc/include/clc/math/clc_half_log2.h +++ b/libclc/clc/include/clc/math/clc_half_log2.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_HALF_LOG2_H__ diff --git a/libclc/clc/include/clc/math/clc_half_powr.h b/libclc/clc/include/clc/math/clc_half_powr.h index 252bfcd785043..bde0c0249abfd 100644 --- a/libclc/clc/include/clc/math/clc_half_powr.h +++ b/libclc/clc/include/clc/math/clc_half_powr.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_HALF_POWR_H__ diff --git a/libclc/clc/include/clc/math/clc_half_recip.h b/libclc/clc/include/clc/math/clc_half_recip.h index 1adb0bbe249cc..40793cfb2aef5 100644 --- a/libclc/clc/include/clc/math/clc_half_recip.h +++ b/libclc/clc/include/clc/math/clc_half_recip.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_HALF_RECIP_H__ diff --git a/libclc/clc/include/clc/math/clc_half_rsqrt.h b/libclc/clc/include/clc/math/clc_half_rsqrt.h index 6342739ed6f20..c0c9971480f87 100644 --- a/libclc/clc/include/clc/math/clc_half_rsqrt.h +++ b/libclc/clc/include/clc/math/clc_half_rsqrt.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_HALF_RSQRT_H__ diff --git a/libclc/clc/include/clc/math/clc_half_sin.h b/libclc/clc/include/clc/math/clc_half_sin.h index 31378d290fecc..022b6bbeabf16 100644 --- a/libclc/clc/include/clc/math/clc_half_sin.h +++ b/libclc/clc/include/clc/math/clc_half_sin.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_HALF_SIN_H__ diff --git a/libclc/clc/include/clc/math/clc_half_sqrt.h b/libclc/clc/include/clc/math/clc_half_sqrt.h index 0e765fb3e7731..be663c298b89c 100644 --- a/libclc/clc/include/clc/math/clc_half_sqrt.h +++ b/libclc/clc/include/clc/math/clc_half_sqrt.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_HALF_SQRT_H__ diff --git a/libclc/clc/include/clc/math/clc_half_tan.h b/libclc/clc/include/clc/math/clc_half_tan.h index 6c890a952522c..acb30c4e31d4e 100644 --- a/libclc/clc/include/clc/math/clc_half_tan.h +++ b/libclc/clc/include/clc/math/clc_half_tan.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_HALF_TAN_H__ diff --git a/libclc/clc/include/clc/math/clc_native_cos.h b/libclc/clc/include/clc/math/clc_native_cos.h index f5837d5479133..bd678301fa505 100644 --- a/libclc/clc/include/clc/math/clc_native_cos.h +++ b/libclc/clc/include/clc/math/clc_native_cos.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_NATIVE_COS_H__ diff --git a/libclc/clc/include/clc/math/clc_native_divide.h b/libclc/clc/include/clc/math/clc_native_divide.h index ed8118775e98b..1e912152cbb6c 100644 --- a/libclc/clc/include/clc/math/clc_native_divide.h +++ b/libclc/clc/include/clc/math/clc_native_divide.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_NATIVE_DIVIDE_H__ diff --git a/libclc/clc/include/clc/math/clc_native_exp.h b/libclc/clc/include/clc/math/clc_native_exp.h index b9c4005585daa..e85d7a8b006e3 100644 --- a/libclc/clc/include/clc/math/clc_native_exp.h +++ b/libclc/clc/include/clc/math/clc_native_exp.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_NATIVE_EXP_H__ diff --git a/libclc/clc/include/clc/math/clc_native_exp10.h b/libclc/clc/include/clc/math/clc_native_exp10.h index 5c49e46330d60..17213bbd2b248 100644 --- a/libclc/clc/include/clc/math/clc_native_exp10.h +++ b/libclc/clc/include/clc/math/clc_native_exp10.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_NATIVE_EXP10_H__ diff --git a/libclc/clc/include/clc/math/clc_native_exp2.h b/libclc/clc/include/clc/math/clc_native_exp2.h index 2fa368f69fef4..2310fcd60cacf 100644 --- a/libclc/clc/include/clc/math/clc_native_exp2.h +++ b/libclc/clc/include/clc/math/clc_native_exp2.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_NATIVE_EXP2_H__ diff --git a/libclc/clc/include/clc/math/clc_native_log.h b/libclc/clc/include/clc/math/clc_native_log.h index 61fe0c0902d85..0613e34f500c8 100644 --- a/libclc/clc/include/clc/math/clc_native_log.h +++ b/libclc/clc/include/clc/math/clc_native_log.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_NATIVE_LOG_H__ diff --git a/libclc/clc/include/clc/math/clc_native_log10.h b/libclc/clc/include/clc/math/clc_native_log10.h index f9fca1d94aedb..462226393c093 100644 --- a/libclc/clc/include/clc/math/clc_native_log10.h +++ b/libclc/clc/include/clc/math/clc_native_log10.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_NATIVE_LOG10_H__ diff --git a/libclc/clc/include/clc/math/clc_native_log2.h b/libclc/clc/include/clc/math/clc_native_log2.h index 69aec752b832b..357ed117097f4 100644 --- a/libclc/clc/include/clc/math/clc_native_log2.h +++ b/libclc/clc/include/clc/math/clc_native_log2.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_NATIVE_LOG2_H__ diff --git a/libclc/clc/include/clc/math/clc_native_powr.h b/libclc/clc/include/clc/math/clc_native_powr.h index 8d50f21a3124d..378b9c8895ea9 100644 --- a/libclc/clc/include/clc/math/clc_native_powr.h +++ b/libclc/clc/include/clc/math/clc_native_powr.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_NATIVE_POWR_H__ diff --git a/libclc/clc/include/clc/math/clc_native_recip.h b/libclc/clc/include/clc/math/clc_native_recip.h index 94d494ad07a17..8ff18012c5f02 100644 --- a/libclc/clc/include/clc/math/clc_native_recip.h +++ b/libclc/clc/include/clc/math/clc_native_recip.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_NATIVE_RECIP_H__ diff --git a/libclc/clc/include/clc/math/clc_native_rsqrt.h b/libclc/clc/include/clc/math/clc_native_rsqrt.h index 5b58d6efd7584..91e3b87261c27 100644 --- a/libclc/clc/include/clc/math/clc_native_rsqrt.h +++ b/libclc/clc/include/clc/math/clc_native_rsqrt.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_NATIVE_RSQRT_H__ diff --git a/libclc/clc/include/clc/math/clc_native_sin.h b/libclc/clc/include/clc/math/clc_native_sin.h index 27f5ecbf40e72..a2354cee90f27 100644 --- a/libclc/clc/include/clc/math/clc_native_sin.h +++ b/libclc/clc/include/clc/math/clc_native_sin.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_NATIVE_SIN_H__ diff --git a/libclc/clc/include/clc/math/clc_native_sqrt.h b/libclc/clc/include/clc/math/clc_native_sqrt.h index a45f8b7f21a21..94fdfd955cedc 100644 --- a/libclc/clc/include/clc/math/clc_native_sqrt.h +++ b/libclc/clc/include/clc/math/clc_native_sqrt.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_NATIVE_SQRT_H__ diff --git a/libclc/clc/include/clc/math/clc_native_tan.h b/libclc/clc/include/clc/math/clc_native_tan.h index 362088a5d58f3..3d300bb3cbcb1 100644 --- a/libclc/clc/include/clc/math/clc_native_tan.h +++ b/libclc/clc/include/clc/math/clc_native_tan.h @@ -16,6 +16,5 @@ #include <clc/math/gentype.inc> #undef __CLC_FUNCTION -#undef __FLOAT_ONLY #endif // __CLC_MATH_CLC_NATIVE_TAN_H__ diff --git a/libclc/clc/include/clc/math/clc_sincos_helpers.h b/libclc/clc/include/clc/math/clc_sincos_helpers.h index e3a2e1c0b605c..f9ceba3bf2cf7 100644 --- a/libclc/clc/include/clc/math/clc_sincos_helpers.h +++ b/libclc/clc/include/clc/math/clc_sincos_helpers.h @@ -14,13 +14,9 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY - #define __DOUBLE_ONLY #define __CLC_BODY <clc/math/clc_sincos_helpers_fp64.inc> #include <clc/math/gentype.inc> -#undef __DOUBLE_ONLY - #endif // __CLC_MATH_CLC_SINCOS_HELPERS_H__ diff --git a/libclc/clc/include/clc/math/gentype.inc b/libclc/clc/include/clc/math/gentype.inc index a8d8317a9cacb..d74a9932bd0a6 100644 --- a/libclc/clc/include/clc/math/gentype.inc +++ b/libclc/clc/include/clc/math/gentype.inc @@ -349,3 +349,7 @@ #undef __CLC_AS_GENTYPE #undef __CLC_CONVERT_GENTYPE + +#undef __HALF_ONLY +#undef __FLOAT_ONLY +#undef __DOUBLE_ONLY diff --git a/libclc/clc/include/clc/shared/binary_def_scalarize.inc b/libclc/clc/include/clc/shared/binary_def_scalarize.inc new file mode 100644 index 0000000000000..3e58e2edacc15 --- /dev/null +++ b/libclc/clc/include/clc/shared/binary_def_scalarize.inc @@ -0,0 +1,121 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include <clc/utils.h> + +#if __CLC_VECSIZE_OR_1 == 1 + +#ifndef __IMPL_FUNCTION +#define __IMPL_FUNCTION FUNCTION +#endif + +#ifndef __CLC_DEF_SPEC +#define __CLC_DEF_SPEC _CLC_DEF +#endif + +#ifndef __CLC_RET_TYPE +#define __CLC_RET_TYPE __CLC_GENTYPE +#endif + +#ifndef __CLC_ARG1_TYPE +#define __CLC_ARG1_TYPE __CLC_GENTYPE +#endif + +#ifndef __CLC_ARG2_TYPE +#define __CLC_ARG2_TYPE __CLC_GENTYPE +#endif + +#ifdef __CLC_HAS_SCALAR +_CLC_OVERLOAD __CLC_DEF_SPEC __CLC_RET_TYPE FUNCTION(__CLC_ARG1_TYPE x, + __CLC_ARG2_TYPE y) { + return __IMPL_FUNCTION(x, y); +} +#endif + +#define __CLC_RET_TYPE2 __CLC_XCONCAT(__CLC_RET_TYPE, 2) +#define __CLC_ARG1_TYPE2 __CLC_XCONCAT(__CLC_ARG1_TYPE, 2) +#define __CLC_ARG2_TYPE2 __CLC_XCONCAT(__CLC_ARG2_TYPE, 2) +_CLC_OVERLOAD __CLC_DEF_SPEC __CLC_RET_TYPE2 FUNCTION(__CLC_ARG1_TYPE2 x, + __CLC_ARG2_TYPE2 y) { + return (__CLC_RET_TYPE2)(__IMPL_FUNCTION(x.s0, y.s0), + __IMPL_FUNCTION(x.s1, y.s1)); +} +#undef __CLC_RET_TYPE2 +#undef __CLC_ARG1_TYPE2 +#undef __CLC_ARG2_TYPE2 + +#define __CLC_RET_TYPE3 __CLC_XCONCAT(__CLC_RET_TYPE, 3) +#define __CLC_ARG1_TYPE3 __CLC_XCONCAT(__CLC_ARG1_TYPE, 3) +#define __CLC_ARG2_TYPE3 __CLC_XCONCAT(__CLC_ARG2_TYPE, 3) +_CLC_OVERLOAD __CLC_DEF_SPEC __CLC_RET_TYPE3 FUNCTION(__CLC_ARG1_TYPE3 x, + __CLC_ARG2_TYPE3 y) { + return (__CLC_RET_TYPE3)(__IMPL_FUNCTION(x.s0, y.s0), + __IMPL_FUNCTION(x.s1, y.s1), + __IMPL_FUNCTION(x.s2, y.s2)); +} +#undef __CLC_RET_TYPE3 +#undef __CLC_ARG1_TYPE3 +#undef __CLC_ARG2_TYPE3 + +#define __CLC_RET_TYPE4 __CLC_XCONCAT(__CLC_RET_TYPE, 4) +#define __CLC_ARG1_TYPE4 __CLC_XCONCAT(__CLC_ARG1_TYPE, 4) +#define __CLC_ARG2_TYPE4 __CLC_XCONCAT(__CLC_ARG2_TYPE, 4) +_CLC_OVERLOAD __CLC_DEF_SPEC __CLC_RET_TYPE4 FUNCTION(__CLC_ARG1_TYPE4 x, + __CLC_ARG2_TYPE4 y) { + return (__CLC_RET_TYPE4)(__IMPL_FUNCTION(x.s0, y.s0), + __IMPL_FUNCTION(x.s1, y.s1), + __IMPL_FUNCTION(x.s2, y.s2), + __IMPL_FUNCTION(x.s3, y.s3)); +} +#undef __CLC_RET_TYPE4 +#undef __CLC_ARG1_TYPE4 +#undef __CLC_ARG2_TYPE4 + +#define __CLC_RET_TYPE8 __CLC_XCONCAT(__CLC_RET_TYPE, 8) +#define __CLC_ARG1_TYPE8 __CLC_XCONCAT(__CLC_ARG1_TYPE, 8) +#define __CLC_ARG2_TYPE8 __CLC_XCONCAT(__CLC_ARG2_TYPE, 8) +_CLC_OVERLOAD __CLC_DEF_SPEC __CLC_RET_TYPE8 FUNCTION(__CLC_ARG1_TYPE8 x, + __CLC_ARG2_TYPE8 y) { + return ( + __CLC_RET_TYPE8)(__IMPL_FUNCTION(x.s0, y.s0), __IMPL_FUNCTION(x.s1, y.s1), + __IMPL_FUNCTION(x.s2, y.s2), __IMPL_FUNCTION(x.s3, y.s3), + __IMPL_FUNCTION(x.s4, y.s4), __IMPL_FUNCTION(x.s5, y.s5), + __IMPL_FUNCTION(x.s6, y.s6), + __IMPL_FUNCTION(x.s7, y.s7)); +} +#undef __CLC_RET_TYPE8 +#undef __CLC_ARG1_TYPE8 +#undef __CLC_ARG2_TYPE8 + +#define __CLC_RET_TYPE16 __CLC_XCONCAT(__CLC_RET_TYPE, 16) +#define __CLC_ARG1_TYPE16 __CLC_XCONCAT(__CLC_ARG1_TYPE, 16) +#define __CLC_ARG2_TYPE16 __CLC_XCONCAT(__CLC_ARG2_TYPE, 16) +_CLC_OVERLOAD __CLC_DEF_SPEC __CLC_RET_TYPE16 FUNCTION(__CLC_ARG1_TYPE16 x, + __CLC_ARG2_TYPE16 y) { + return (__CLC_RET_TYPE16)(__IMPL_FUNCTION(x.s0, y.s0), + __IMPL_FUNCTION(x.s1, y.s1), + __IMPL_FUNCTION(x.s2, y.s2), + __IMPL_FUNCTION(x.s3, y.s3), + __IMPL_FUNCTION(x.s4, y.s4), + __IMPL_FUNCTION(x.s5, y.s5), + __IMPL_FUNCTION(x.s6, y.s6), + __IMPL_FUNCTION(x.s7, y.s7), + __IMPL_FUNCTION(x.s8, y.s8), + __IMPL_FUNCTION(x.s9, y.s9), + __IMPL_FUNCTION(x.sa, y.sa), + __IMPL_FUNCTION(x.sb, y.sb), + __IMPL_FUNCTION(x.sc, y.sc), + __IMPL_FUNCTION(x.sd, y.sd), + __IMPL_FUNCTION(x.se, y.se), + __IMPL_FUNCTION(x.sf, y.sf)); +} +#undef __CLC_RET_TYPE16 +#undef __CLC_ARG1_TYPE16 +#undef __CLC_ARG2_TYPE16 + +#endif // __CLC_VECSIZE_OR_1 == 1 diff --git a/libclc/clc/include/clc/shared/ternary_def_scalarize.inc b/libclc/clc/include/clc/shared/ternary_def_scalarize.inc new file mode 100644 index 0000000000000..7a66e34a9229f --- /dev/null +++ b/libclc/clc/include/clc/shared/ternary_def_scalarize.inc @@ -0,0 +1,143 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include <clc/utils.h> + +#if __CLC_VECSIZE_OR_1 == 1 + +#ifndef __IMPL_FUNCTION +#define __IMPL_FUNCTION FUNCTION +#endif + +#ifndef __CLC_DEF_SPEC +#define __CLC_DEF_SPEC _CLC_DEF +#endif + +#ifndef __CLC_RET_TYPE +#define __CLC_RET_TYPE __CLC_GENTYPE +#endif + +#ifndef __CLC_ARG1_TYPE +#define __CLC_ARG1_TYPE __CLC_GENTYPE +#endif + +#ifndef __CLC_ARG2_TYPE +#define __CLC_ARG2_TYPE __CLC_GENTYPE +#endif + +#ifndef __CLC_ARG3_TYPE +#define __CLC_ARG3_TYPE __CLC_GENTYPE +#endif + +#ifdef __CLC_HAS_SCALAR +_CLC_OVERLOAD __CLC_DEF_SPEC __CLC_RET_TYPE FUNCTION(__CLC_ARG1_TYPE x, + __CLC_ARG2_TYPE y, + __CLC_ARG3_TYPE z) { + return __IMPL_FUNCTION(x, y, z); +} +#endif + +#define __CLC_RET_TYPE2 __CLC_XCONCAT(__CLC_RET_TYPE, 2) +#define __CLC_ARG1_TYPE2 __CLC_XCONCAT(__CLC_ARG1_TYPE, 2) +#define __CLC_ARG2_TYPE2 __CLC_XCONCAT(__CLC_ARG2_TYPE, 2) +#define __CLC_ARG3_TYPE2 __CLC_XCONCAT(__CLC_ARG3_TYPE, 2) +_CLC_OVERLOAD __CLC_DEF_SPEC __CLC_RET_TYPE2 FUNCTION(__CLC_ARG1_TYPE2 x, + __CLC_ARG2_TYPE2 y, + __CLC_ARG3_TYPE2 z) { + return (__CLC_RET_TYPE2)(__IMPL_FUNCTION(x.s0, y.s0, z.s0), + __IMPL_FUNCTION(x.s1, y.s1, z.s1)); +} +#undef __CLC_RET_TYPE2 +#undef __CLC_ARG1_TYPE2 +#undef __CLC_ARG2_TYPE2 +#undef __CLC_ARG3_TYPE2 + +#define __CLC_RET_TYPE3 __CLC_XCONCAT(__CLC_RET_TYPE, 3) +#define __CLC_ARG1_TYPE3 __CLC_XCONCAT(__CLC_ARG1_TYPE, 3) +#define __CLC_ARG2_TYPE3 __CLC_XCONCAT(__CLC_ARG2_TYPE, 3) +#define __CLC_ARG3_TYPE3 __CLC_XCONCAT(__CLC_ARG3_TYPE, 3) +_CLC_OVERLOAD __CLC_DEF_SPEC __CLC_RET_TYPE3 FUNCTION(__CLC_ARG1_TYPE3 x, + __CLC_ARG2_TYPE3 y, + __CLC_ARG3_TYPE3 z) { + return (__CLC_RET_TYPE3)(__IMPL_FUNCTION(x.s0, y.s0, z.s0), + __IMPL_FUNCTION(x.s1, y.s1, z.s1), + __IMPL_FUNCTION(x.s2, y.s2, z.s2)); +} +#undef __CLC_RET_TYPE3 +#undef __CLC_ARG1_TYPE3 +#undef __CLC_ARG2_TYPE3 +#undef __CLC_ARG3_TYPE3 + +#define __CLC_RET_TYPE4 __CLC_XCONCAT(__CLC_RET_TYPE, 4) +#define __CLC_ARG1_TYPE4 __CLC_XCONCAT(__CLC_ARG1_TYPE, 4) +#define __CLC_ARG2_TYPE4 __CLC_XCONCAT(__CLC_ARG2_TYPE, 4) +#define __CLC_ARG3_TYPE4 __CLC_XCONCAT(__CLC_ARG3_TYPE, 4) +_CLC_OVERLOAD __CLC_DEF_SPEC __CLC_RET_TYPE4 FUNCTION(__CLC_ARG1_TYPE4 x, + __CLC_ARG2_TYPE4 y, + __CLC_ARG3_TYPE4 z) { + return (__CLC_RET_TYPE4)(__IMPL_FUNCTION(x.s0, y.s0, z.s0), + __IMPL_FUNCTION(x.s1, y.s1, z.s1), + __IMPL_FUNCTION(x.s2, y.s2, z.s2), + __IMPL_FUNCTION(x.s3, y.s3, z.s3)); +} +#undef __CLC_RET_TYPE4 +#undef __CLC_ARG1_TYPE4 +#undef __CLC_ARG2_TYPE4 +#undef __CLC_ARG3_TYPE4 + +#define __CLC_RET_TYPE8 __CLC_XCONCAT(__CLC_RET_TYPE, 8) +#define __CLC_ARG1_TYPE8 __CLC_XCONCAT(__CLC_ARG1_TYPE, 8) +#define __CLC_ARG2_TYPE8 __CLC_XCONCAT(__CLC_ARG2_TYPE, 8) +#define __CLC_ARG3_TYPE8 __CLC_XCONCAT(__CLC_ARG3_TYPE, 8) +_CLC_OVERLOAD __CLC_DEF_SPEC __CLC_RET_TYPE8 FUNCTION(__CLC_ARG1_TYPE8 x, + __CLC_ARG2_TYPE8 y, + __CLC_ARG3_TYPE8 z) { + return (__CLC_RET_TYPE8)(__IMPL_FUNCTION(x.s0, y.s0, z.s0), + __IMPL_FUNCTION(x.s1, y.s1, z.s1), + __IMPL_FUNCTION(x.s2, y.s2, z.s2), + __IMPL_FUNCTION(x.s3, y.s3, z.s3), + __IMPL_FUNCTION(x.s4, y.s4, z.s4), + __IMPL_FUNCTION(x.s5, y.s5, z.s5), + __IMPL_FUNCTION(x.s6, y.s6, z.s6), + __IMPL_FUNCTION(x.s7, y.s7, z.s7)); +} +#undef __CLC_RET_TYPE8 +#undef __CLC_ARG1_TYPE8 +#undef __CLC_ARG2_TYPE8 +#undef __CLC_ARG3_TYPE8 + +#define __CLC_RET_TYPE16 __CLC_XCONCAT(__CLC_RET_TYPE, 16) +#define __CLC_ARG1_TYPE16 __CLC_XCONCAT(__CLC_ARG1_TYPE, 16) +#define __CLC_ARG2_TYPE16 __CLC_XCONCAT(__CLC_ARG2_TYPE, 16) +#define __CLC_ARG3_TYPE16 __CLC_XCONCAT(__CLC_ARG3_TYPE, 16) +_CLC_OVERLOAD __CLC_DEF_SPEC __CLC_RET_TYPE16 FUNCTION(__CLC_ARG1_TYPE16 x, + __CLC_ARG2_TYPE16 y, + __CLC_ARG3_TYPE16 z) { + return (__CLC_RET_TYPE16)(__IMPL_FUNCTION(x.s0, y.s0, z.s0), + __IMPL_FUNCTION(x.s1, y.s1, z.s1), + __IMPL_FUNCTION(x.s2, y.s2, z.s2), + __IMPL_FUNCTION(x.s3, y.s3, z.s3), + __IMPL_FUNCTION(x.s4, y.s4, z.s4), + __IMPL_FUNCTION(x.s5, y.s5, z.s5), + __IMPL_FUNCTION(x.s6, y.s6, z.s6), + __IMPL_FUNCTION(x.s7, y.s7, z.s7), + __IMPL_FUNCTION(x.s8, y.s8, z.s8), + __IMPL_FUNCTION(x.s9, y.s9, z.s9), + __IMPL_FUNCTION(x.sa, y.sa, z.sa), + __IMPL_FUNCTION(x.sb, y.sb, z.sb), + __IMPL_FUNCTION(x.sc, y.sc, z.sc), + __IMPL_FUNCTION(x.sd, y.sd, z.sd), + __IMPL_FUNCTION(x.se, y.se, z.se), + __IMPL_FUNCTION(x.sf, y.sf, z.sf)); +} +#undef __CLC_RET_TYPE16 +#undef __CLC_ARG1_TYPE16 +#undef __CLC_ARG2_TYPE16 +#undef __CLC_ARG3_TYPE16 + +#endif // __CLC_VECSIZE_OR_1 == 1 diff --git a/libclc/clc/include/clc/shared/unary_def_scalarize.inc b/libclc/clc/include/clc/shared/unary_def_scalarize.inc new file mode 100644 index 0000000000000..f0d0338441354 --- /dev/null +++ b/libclc/clc/include/clc/shared/unary_def_scalarize.inc @@ -0,0 +1,87 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include <clc/utils.h> + +#if __CLC_VECSIZE_OR_1 == 1 + +#ifndef __IMPL_FUNCTION +#define __IMPL_FUNCTION FUNCTION +#endif + +#ifndef __CLC_RET_TYPE +#define __CLC_RET_TYPE __CLC_GENTYPE +#endif + +#ifndef __CLC_ARG1_TYPE +#define __CLC_ARG1_TYPE __CLC_GENTYPE +#endif + +#ifndef __CLC_ARG2_TYPE +#define __CLC_ARG2_TYPE __CLC_GENTYPE +#endif + +#ifdef __CLC_HAS_SCALAR +_CLC_OVERLOAD _CLC_DEF __CLC_RET_TYPE FUNCTION(__CLC_ARG1_TYPE x) { + return __IMPL_FUNCTION(x); +} +#endif + +#define __CLC_RET_TYPE2 __CLC_XCONCAT(__CLC_RET_TYPE, 2) +#define __CLC_ARG1_TYPE2 __CLC_XCONCAT(__CLC_ARG1_TYPE, 2) +_CLC_OVERLOAD _CLC_DEF __CLC_RET_TYPE2 FUNCTION(__CLC_ARG1_TYPE2 x) { + return (__CLC_RET_TYPE2)(__IMPL_FUNCTION(x.s0), __IMPL_FUNCTION(x.s1)); +} +#undef __CLC_RET_TYPE2 +#undef __CLC_ARG1_TYPE2 + +#define __CLC_RET_TYPE3 __CLC_XCONCAT(__CLC_RET_TYPE, 3) +#define __CLC_ARG1_TYPE3 __CLC_XCONCAT(__CLC_ARG1_TYPE, 3) +_CLC_OVERLOAD _CLC_DEF __CLC_RET_TYPE3 FUNCTION(__CLC_ARG1_TYPE3 x) { + return (__CLC_RET_TYPE3)(__IMPL_FUNCTION(x.s0), __IMPL_FUNCTION(x.s1), + __IMPL_FUNCTION(x.s2)); +} +#undef __CLC_RET_TYPE3 +#undef __CLC_ARG1_TYPE3 + +#define __CLC_RET_TYPE4 __CLC_XCONCAT(__CLC_RET_TYPE, 4) +#define __CLC_ARG1_TYPE4 __CLC_XCONCAT(__CLC_ARG1_TYPE, 4) +_CLC_OVERLOAD _CLC_DEF __CLC_RET_TYPE4 FUNCTION(__CLC_ARG1_TYPE4 x) { + return (__CLC_RET_TYPE4)(__IMPL_FUNCTION(x.s0), __IMPL_FUNCTION(x.s1), + __IMPL_FUNCTION(x.s2), __IMPL_FUNCTION(x.s3)); +} +#undef __CLC_RET_TYPE4 +#undef __CLC_ARG1_TYPE4 + +#define __CLC_RET_TYPE8 __CLC_XCONCAT(__CLC_RET_TYPE, 8) +#define __CLC_ARG1_TYPE8 __CLC_XCONCAT(__CLC_ARG1_TYPE, 8) +_CLC_OVERLOAD _CLC_DEF __CLC_RET_TYPE8 FUNCTION(__CLC_ARG1_TYPE8 x) { + return (__CLC_RET_TYPE8)(__IMPL_FUNCTION(x.s0), __IMPL_FUNCTION(x.s1), + __IMPL_FUNCTION(x.s2), __IMPL_FUNCTION(x.s3), + __IMPL_FUNCTION(x.s4), __IMPL_FUNCTION(x.s5), + __IMPL_FUNCTION(x.s6), __IMPL_FUNCTION(x.s7)); +} +#undef __CLC_RET_TYPE8 +#undef __CLC_ARG1_TYPE8 + +#define __CLC_RET_TYPE16 __CLC_XCONCAT(__CLC_RET_TYPE, 16) +#define __CLC_ARG1_TYPE16 __CLC_XCONCAT(__CLC_ARG1_TYPE, 16) +_CLC_OVERLOAD _CLC_DEF __CLC_RET_TYPE16 FUNCTION(__CLC_ARG1_TYPE16 x) { + return (__CLC_RET_TYPE16)(__IMPL_FUNCTION(x.s0), __IMPL_FUNCTION(x.s1), + __IMPL_FUNCTION(x.s2), __IMPL_FUNCTION(x.s3), + __IMPL_FUNCTION(x.s4), __IMPL_FUNCTION(x.s5), + __IMPL_FUNCTION(x.s6), __IMPL_FUNCTION(x.s7), + __IMPL_FUNCTION(x.s8), __IMPL_FUNCTION(x.s9), + __IMPL_FUNCTION(x.sa), __IMPL_FUNCTION(x.sb), + __IMPL_FUNCTION(x.sc), __IMPL_FUNCTION(x.sd), + __IMPL_FUNCTION(x.se), __IMPL_FUNCTION(x.sf)); +} +#undef __CLC_RET_TYPE16 +#undef __CLC_ARG1_TYPE16 + +#endif // __CLC_VECSIZE_OR_1 == 1 diff --git a/libclc/clc/lib/amdgcn/math/clc_fmax.cl b/libclc/clc/lib/amdgcn/math/clc_fmax.cl index 20bdcadb9eabf..cea90a7135d5a 100644 --- a/libclc/clc/lib/amdgcn/math/clc_fmax.cl +++ b/libclc/clc/lib/amdgcn/math/clc_fmax.cl @@ -18,7 +18,6 @@ _CLC_DEF _CLC_OVERLOAD float __clc_fmax(float x, float y) { y = __builtin_canonicalizef(y); return __builtin_fmaxf(x, y); } -_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_fmax, float, float) #ifdef cl_khr_fp64 @@ -29,8 +28,6 @@ _CLC_DEF _CLC_OVERLOAD double __clc_fmax(double x, double y) { y = __builtin_canonicalize(y); return __builtin_fmax(x, y); } -_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, __clc_fmax, double, - double) #endif #ifdef cl_khr_fp16 @@ -44,6 +41,9 @@ _CLC_DEF _CLC_OVERLOAD half __clc_fmax(half x, half y) { return x; return (y < x) ? x : y; } -_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, half, __clc_fmax, half, half) #endif + +#define FUNCTION __clc_fmax +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> diff --git a/libclc/clc/lib/amdgcn/math/clc_fmin.cl b/libclc/clc/lib/amdgcn/math/clc_fmin.cl index a5f66dfefa900..12bb0c64429fd 100644 --- a/libclc/clc/lib/amdgcn/math/clc_fmin.cl +++ b/libclc/clc/lib/amdgcn/math/clc_fmin.cl @@ -18,7 +18,6 @@ _CLC_DEF _CLC_OVERLOAD float __clc_fmin(float x, float y) { y = __builtin_canonicalizef(y); return __builtin_fminf(x, y); } -_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_fmin, float, float) #ifdef cl_khr_fp64 @@ -29,8 +28,6 @@ _CLC_DEF _CLC_OVERLOAD double __clc_fmin(double x, double y) { y = __builtin_canonicalize(y); return __builtin_fmin(x, y); } -_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, __clc_fmin, double, - double) #endif @@ -45,6 +42,9 @@ _CLC_DEF _CLC_OVERLOAD half __clc_fmin(half x, half y) { return x; return (y < x) ? y : x; } -_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, half, __clc_fmin, half, half) #endif + +#define FUNCTION __clc_fmin +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> diff --git a/libclc/clc/lib/amdgcn/math/clc_ldexp_override.cl b/libclc/clc/lib/amdgcn/math/clc_ldexp_override.cl index 94a6d6ccfa9f8..150496ade0622 100644 --- a/libclc/clc/lib/amdgcn/math/clc_ldexp_override.cl +++ b/libclc/clc/lib/amdgcn/math/clc_ldexp_override.cl @@ -10,14 +10,25 @@ #include <clc/internal/clc.h> #include <clc/math/clc_ldexp.h> +#define FUNCTION __clc_ldexp +#define __CLC_ARG2_TYPE int +#define __CLC_HAS_SCALAR + #ifdef __HAS_LDEXPF__ // This defines all the ldexp(floatN, intN) variants. -_CLC_DEFINE_BINARY_BUILTIN(float, __clc_ldexp, __builtin_amdgcn_ldexpf, float, int); +#define __FLOAT_ONLY +#define __IMPL_FUNCTION __builtin_amdgcn_ldexpf +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef __IMPL_FUNCTION #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, __clc_ldexp, __builtin_amdgcn_ldexp, double, - int); +#define __DOUBLE_ONLY +#define __IMPL_FUNCTION __builtin_amdgcn_ldexp +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef __IMPL_FUNCTION #endif diff --git a/libclc/clc/lib/amdgpu/math/clc_native_exp2.cl b/libclc/clc/lib/amdgpu/math/clc_native_exp2.cl index 76b1850fce574..9799fefeb0a01 100644 --- a/libclc/clc/lib/amdgpu/math/clc_native_exp2.cl +++ b/libclc/clc/lib/amdgpu/math/clc_native_exp2.cl @@ -9,8 +9,9 @@ #include <clc/clcmacro.h> #include <clc/internal/clc.h> -_CLC_OVERLOAD _CLC_DEF float __clc_native_exp2(float val) { - return __builtin_amdgcn_exp2f(val); -} - -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_native_exp2, float) +#define __FLOAT_ONLY +#define __CLC_HAS_SCALAR +#define FUNCTION __clc_native_exp2 +#define __IMPL_FUNCTION __builtin_amdgcn_exp2f +#define __CLC_BODY <clc/shared/unary_def_scalarize.inc> +#include <clc/math/gentype.inc> diff --git a/libclc/clc/lib/amdgpu/math/clc_sqrt_fp64.cl b/libclc/clc/lib/amdgpu/math/clc_sqrt_fp64.cl index 576525bc47f6a..b7cb635a2ae8e 100644 --- a/libclc/clc/lib/amdgpu/math/clc_sqrt_fp64.cl +++ b/libclc/clc/lib/amdgpu/math/clc_sqrt_fp64.cl @@ -43,6 +43,9 @@ _CLC_OVERLOAD _CLC_DEF double __clc_sqrt(double x) { return (x == __builtin_inf() || (x == 0.0)) ? v01 : v23; } -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, __clc_sqrt, double); +#define __DOUBLE_ONLY +#define FUNCTION __clc_sqrt +#define __CLC_BODY <clc/shared/unary_def_scalarize.inc> +#include <clc/math/gentype.inc> #endif diff --git a/libclc/clc/lib/clspv/math/clc_sw_fma.cl b/libclc/clc/lib/clspv/math/clc_sw_fma.cl index 82978380a63f6..266269644721a 100644 --- a/libclc/clc/lib/clspv/math/clc_sw_fma.cl +++ b/libclc/clc/lib/clspv/math/clc_sw_fma.cl @@ -269,5 +269,7 @@ _CLC_DEF _CLC_OVERLOAD float __clc_sw_fma(float a, float b, float c) { ((uint)st_fma.mantissa.lo & 0x7fffff)); } -_CLC_TERNARY_VECTORIZE(_CLC_DEF _CLC_OVERLOAD, float, __clc_sw_fma, float, - float, float) +#define __FLOAT_ONLY +#define FUNCTION __clc_sw_fma +#define __CLC_BODY <clc/shared/ternary_def_scalarize.inc> +#include <clc/math/gentype.inc> diff --git a/libclc/clc/lib/generic/integer/clc_clz.cl b/libclc/clc/lib/generic/integer/clc_clz.cl index 74f662375af6b..c6e1da680b7bd 100644 --- a/libclc/clc/lib/generic/integer/clc_clz.cl +++ b/libclc/clc/lib/generic/integer/clc_clz.cl @@ -42,11 +42,6 @@ _CLC_OVERLOAD _CLC_DEF ulong __clc_clz(ulong x) { return x ? __builtin_clzl(x) : 64; } -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, char, __clc_clz, char) -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, uchar, __clc_clz, uchar) -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, short, __clc_clz, short) -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, ushort, __clc_clz, ushort) -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, int, __clc_clz, int) -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, uint, __clc_clz, uint) -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, long, __clc_clz, long) -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, ulong, __clc_clz, ulong) +#define FUNCTION __clc_clz +#define __CLC_BODY <clc/shared/unary_def_scalarize.inc> +#include <clc/integer/gentype.inc> diff --git a/libclc/clc/lib/generic/integer/clc_ctz.cl b/libclc/clc/lib/generic/integer/clc_ctz.cl index 50fda4a214b24..d82d99d539df8 100644 --- a/libclc/clc/lib/generic/integer/clc_ctz.cl +++ b/libclc/clc/lib/generic/integer/clc_ctz.cl @@ -38,11 +38,6 @@ _CLC_OVERLOAD _CLC_DEF ulong __clc_ctz(ulong x) { return __builtin_ctzg(x, 64); } -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, char, __clc_ctz, char) -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, uchar, __clc_ctz, uchar) -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, short, __clc_ctz, short) -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, ushort, __clc_ctz, ushort) -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, int, __clc_ctz, int) -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, uint, __clc_ctz, uint) -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, long, __clc_ctz, long) -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, ulong, __clc_ctz, ulong) +#define FUNCTION __clc_ctz +#define __CLC_BODY <clc/shared/unary_def_scalarize.inc> +#include <clc/integer/gentype.inc> diff --git a/libclc/clc/lib/generic/math/clc_erf.cl b/libclc/clc/lib/generic/math/clc_erf.cl index 3808b9b81411a..910926d27a343 100644 --- a/libclc/clc/lib/generic/math/clc_erf.cl +++ b/libclc/clc/lib/generic/math/clc_erf.cl @@ -211,7 +211,11 @@ _CLC_OVERLOAD _CLC_DEF float __clc_erf(float x) { return ret; } -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_erf, float); +#define __FLOAT_ONLY +#define FUNCTION __clc_erf +#define __CLC_BODY <clc/shared/unary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef FUNCTION #ifdef cl_khr_fp64 @@ -496,7 +500,11 @@ _CLC_OVERLOAD _CLC_DEF double __clc_erf(double y) { return y < 0.0 ? -ret : ret; } -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, __clc_erf, double); +#define __DOUBLE_ONLY +#define FUNCTION __clc_erf +#define __CLC_BODY <clc/shared/unary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef FUNCTION #endif diff --git a/libclc/clc/lib/generic/math/clc_erfc.cl b/libclc/clc/lib/generic/math/clc_erfc.cl index a9edfdbb72f23..105c299366148 100644 --- a/libclc/clc/lib/generic/math/clc_erfc.cl +++ b/libclc/clc/lib/generic/math/clc_erfc.cl @@ -211,7 +211,11 @@ _CLC_OVERLOAD _CLC_DEF float __clc_erfc(float x) { return ret; } -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_erfc, float); +#define __FLOAT_ONLY +#define FUNCTION __clc_erfc +#define __CLC_BODY <clc/shared/unary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef FUNCTION #ifdef cl_khr_fp64 @@ -505,7 +509,11 @@ _CLC_OVERLOAD _CLC_DEF double __clc_erfc(double x) { return ret; } -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, __clc_erfc, double); +#define __DOUBLE_ONLY +#define FUNCTION __clc_erfc +#define __CLC_BODY <clc/shared/unary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef FUNCTION #endif diff --git a/libclc/clc/lib/generic/math/clc_fmax.cl b/libclc/clc/lib/generic/math/clc_fmax.cl index 8ee369f57d38b..cc8f423daf418 100644 --- a/libclc/clc/lib/generic/math/clc_fmax.cl +++ b/libclc/clc/lib/generic/math/clc_fmax.cl @@ -10,13 +10,29 @@ #include <clc/internal/clc.h> #include <clc/relational/clc_isnan.h> -_CLC_DEFINE_BINARY_BUILTIN(float, __clc_fmax, __builtin_fmaxf, float, float); +#define __FLOAT_ONLY +#define __CLC_HAS_SCALAR +#define FUNCTION __clc_fmax +#define __IMPL_FUNCTION __builtin_fmaxf +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef __CLC_HAS_SCALAR +#undef FUNCTION +#undef __IMPL_FUNCTION #ifdef cl_khr_fp64 #pragma OPENCL EXTENSION cl_khr_fp64 : enable -_CLC_DEFINE_BINARY_BUILTIN(double, __clc_fmax, __builtin_fmax, double, double); +#define __DOUBLE_ONLY +#define __CLC_HAS_SCALAR +#define FUNCTION __clc_fmax +#define __IMPL_FUNCTION __builtin_fmax +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef __CLC_HAS_SCALAR +#undef FUNCTION +#undef __IMPL_FUNCTION #endif @@ -31,6 +47,12 @@ _CLC_DEF _CLC_OVERLOAD half __clc_fmax(half x, half y) { return x; return (x < y) ? y : x; } -_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, half, __clc_fmax, half, half) + +#define __HALF_ONLY +#define __CLC_SUPPORTED_VECSIZE_OR_1 2 +#define FUNCTION __clc_fmax +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef FUNCTION #endif diff --git a/libclc/clc/lib/generic/math/clc_fmin.cl b/libclc/clc/lib/generic/math/clc_fmin.cl index 2f307274b9be5..ef57cce9ed830 100644 --- a/libclc/clc/lib/generic/math/clc_fmin.cl +++ b/libclc/clc/lib/generic/math/clc_fmin.cl @@ -10,13 +10,29 @@ #include <clc/internal/clc.h> #include <clc/relational/clc_isnan.h> -_CLC_DEFINE_BINARY_BUILTIN(float, __clc_fmin, __builtin_fminf, float, float); +#define __FLOAT_ONLY +#define __CLC_HAS_SCALAR +#define FUNCTION __clc_fmin +#define __IMPL_FUNCTION __builtin_fminf +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef __CLC_HAS_SCALAR +#undef FUNCTION +#undef __IMPL_FUNCTION #ifdef cl_khr_fp64 #pragma OPENCL EXTENSION cl_khr_fp64 : enable -_CLC_DEFINE_BINARY_BUILTIN(double, __clc_fmin, __builtin_fmin, double, double); +#define __DOUBLE_ONLY +#define __CLC_HAS_SCALAR +#define FUNCTION __clc_fmin +#define __IMPL_FUNCTION __builtin_fmin +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef __CLC_HAS_SCALAR +#undef FUNCTION +#undef __IMPL_FUNCTION #endif @@ -31,6 +47,11 @@ _CLC_DEF _CLC_OVERLOAD half __clc_fmin(half x, half y) { return x; return (y < x) ? y : x; } -_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, half, __clc_fmin, half, half) + +#define __HALF_ONLY +#define __CLC_SUPPORTED_VECSIZE_OR_1 2 +#define FUNCTION __clc_fmin +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> #endif diff --git a/libclc/clc/lib/generic/math/clc_fmod.cl b/libclc/clc/lib/generic/math/clc_fmod.cl index 6af84a14f3d1a..30ab1789d7941 100644 --- a/libclc/clc/lib/generic/math/clc_fmod.cl +++ b/libclc/clc/lib/generic/math/clc_fmod.cl @@ -63,7 +63,12 @@ _CLC_DEF _CLC_OVERLOAD float __clc_fmod(float x, float y) { return xr; } -_CLC_BINARY_VECTORIZE(_CLC_DEF _CLC_OVERLOAD, float, __clc_fmod, float, float); + +#define __FLOAT_ONLY +#define FUNCTION __clc_fmod +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef FUNCTION #ifdef cl_khr_fp64 @@ -170,8 +175,13 @@ _CLC_DEF _CLC_OVERLOAD double __clc_fmod(double x, double y) { return ret; } -_CLC_BINARY_VECTORIZE(_CLC_DEF _CLC_OVERLOAD, double, __clc_fmod, double, - double); + +#define __DOUBLE_ONLY +#define FUNCTION __clc_fmod +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef FUNCTION + #endif #ifdef cl_khr_fp16 diff --git a/libclc/clc/lib/generic/math/clc_ldexp.cl b/libclc/clc/lib/generic/math/clc_ldexp.cl index 144a8fc897d99..cb4185d89c729 100644 --- a/libclc/clc/lib/generic/math/clc_ldexp.cl +++ b/libclc/clc/lib/generic/math/clc_ldexp.cl @@ -86,8 +86,6 @@ _CLC_DEF_ldexp _CLC_OVERLOAD float __clc_ldexp(float x, int n) { 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 @@ -121,8 +119,6 @@ _CLC_DEF_ldexp _CLC_OVERLOAD double __clc_ldexp(double x, int n) { return mr; } -_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF_ldexp, double, __clc_ldexp, double, int); - #endif #ifdef cl_khr_fp16 @@ -133,6 +129,10 @@ _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_ldexp, half, __clc_ldexp, half, int); - #endif + +#define FUNCTION __clc_ldexp +#define __CLC_DEF_SPEC _CLC_DEF_ldexp +#define __CLC_ARG2_TYPE int +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> diff --git a/libclc/clc/lib/generic/math/clc_log.cl b/libclc/clc/lib/generic/math/clc_log.cl index 85405faf8b8ca..cf5628f206cad 100644 --- a/libclc/clc/lib/generic/math/clc_log.cl +++ b/libclc/clc/lib/generic/math/clc_log.cl @@ -19,8 +19,6 @@ _CLC_OVERLOAD _CLC_DEF float __clc_log(float x) { return __clc_log2(x) * (1.0f / M_LOG2E_F); } -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_log, float); - #ifdef cl_khr_fp64 #pragma OPENCL EXTENSION cl_khr_fp64 : enable @@ -29,8 +27,6 @@ _CLC_OVERLOAD _CLC_DEF double __clc_log(double x) { return __clc_log2(x) * (1.0 / M_LOG2E); } -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, __clc_log, double); - #endif // cl_khr_fp64 #ifdef cl_khr_fp16 @@ -41,6 +37,8 @@ _CLC_OVERLOAD _CLC_DEF half __clc_log(half x) { return (half)__clc_log2((float)x) * (1.0h / M_LOG2E_H); } -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, half, __clc_log, half); - #endif // cl_khr_fp16 + +#define FUNCTION __clc_log +#define __CLC_BODY <clc/shared/unary_def_scalarize.inc> +#include <clc/math/gentype.inc> diff --git a/libclc/clc/lib/generic/math/clc_log10.cl b/libclc/clc/lib/generic/math/clc_log10.cl index 689ec25b74e5f..f5f0e8cc70837 100644 --- a/libclc/clc/lib/generic/math/clc_log10.cl +++ b/libclc/clc/lib/generic/math/clc_log10.cl @@ -22,12 +22,6 @@ #include "clc_log_base.h" #undef COMPILING_LOG10 -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_log10, float); - -#ifdef cl_khr_fp64 -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, __clc_log10, double); -#endif // cl_khr_fp64 - -#ifdef cl_khr_fp16 -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, half, __clc_log10, half); -#endif // cl_khr_fp16 +#define FUNCTION __clc_log10 +#define __CLC_BODY <clc/shared/unary_def_scalarize.inc> +#include <clc/math/gentype.inc> diff --git a/libclc/clc/lib/generic/math/clc_log2.cl b/libclc/clc/lib/generic/math/clc_log2.cl index ebd981c9d9955..335488af2f3de 100644 --- a/libclc/clc/lib/generic/math/clc_log2.cl +++ b/libclc/clc/lib/generic/math/clc_log2.cl @@ -22,12 +22,6 @@ #include "clc_log_base.h" #undef COMPILING_LOG2 -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_log2, float); - -#ifdef cl_khr_fp64 -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, __clc_log2, double); -#endif // cl_khr_fp64 - -#ifdef cl_khr_fp16 -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, half, __clc_log2, half); -#endif // cl_khr_fp16 +#define FUNCTION __clc_log2 +#define __CLC_BODY <clc/shared/unary_def_scalarize.inc> +#include <clc/math/gentype.inc> diff --git a/libclc/clc/lib/generic/math/clc_remainder.cl b/libclc/clc/lib/generic/math/clc_remainder.cl index c79d271c624e2..1d631f22874d3 100644 --- a/libclc/clc/lib/generic/math/clc_remainder.cl +++ b/libclc/clc/lib/generic/math/clc_remainder.cl @@ -73,8 +73,12 @@ _CLC_DEF _CLC_OVERLOAD float __clc_remainder(float x, float y) { return xr; } -_CLC_BINARY_VECTORIZE(_CLC_DEF _CLC_OVERLOAD, float, __clc_remainder, float, - float); + +#define __FLOAT_ONLY +#define FUNCTION __clc_remainder +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef FUNCTION #ifdef cl_khr_fp64 @@ -207,8 +211,13 @@ _CLC_DEF _CLC_OVERLOAD double __clc_remainder(double x, double y) { return ret; } -_CLC_BINARY_VECTORIZE(_CLC_DEF _CLC_OVERLOAD, double, __clc_remainder, double, - double); + +#define __DOUBLE_ONLY +#define FUNCTION __clc_remainder +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef FUNCTION + #endif #ifdef cl_khr_fp16 diff --git a/libclc/clc/lib/generic/math/clc_sincos_helpers.cl b/libclc/clc/lib/generic/math/clc_sincos_helpers.cl index c1768e39243fa..0ea1195fffa70 100644 --- a/libclc/clc/lib/generic/math/clc_sincos_helpers.cl +++ b/libclc/clc/lib/generic/math/clc_sincos_helpers.cl @@ -32,8 +32,6 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY - #ifdef cl_khr_fp64 #pragma OPENCL EXTENSION cl_khr_fp64 : enable @@ -52,6 +50,4 @@ #include <clc/math/gentype.inc> -#undef __DOUBLE_ONLY - #endif diff --git a/libclc/clc/lib/generic/math/clc_sw_fma.cl b/libclc/clc/lib/generic/math/clc_sw_fma.cl index df5ac9dbd49d7..ee4734078d698 100644 --- a/libclc/clc/lib/generic/math/clc_sw_fma.cl +++ b/libclc/clc/lib/generic/math/clc_sw_fma.cl @@ -160,5 +160,7 @@ _CLC_DEF _CLC_OVERLOAD float __clc_sw_fma(float a, float b, float c) { ((uint)st_fma.mantissa & 0x7fffff)); } -_CLC_TERNARY_VECTORIZE(_CLC_DEF _CLC_OVERLOAD, float, __clc_sw_fma, float, - float, float) +#define __FLOAT_ONLY +#define FUNCTION __clc_sw_fma +#define __CLC_BODY <clc/shared/ternary_def_scalarize.inc> +#include <clc/math/gentype.inc> diff --git a/libclc/clc/lib/generic/math/clc_tgamma.cl b/libclc/clc/lib/generic/math/clc_tgamma.cl index 5c2d164abeee8..83d9942eaf3f7 100644 --- a/libclc/clc/lib/generic/math/clc_tgamma.cl +++ b/libclc/clc/lib/generic/math/clc_tgamma.cl @@ -32,7 +32,11 @@ _CLC_OVERLOAD _CLC_DEF float __clc_tgamma(float x) { return g; } -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_tgamma, float); +#define __FLOAT_ONLY +#define FUNCTION __clc_tgamma +#define __CLC_BODY <clc/shared/unary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef FUNCTION #ifdef cl_khr_fp64 @@ -55,7 +59,11 @@ _CLC_OVERLOAD _CLC_DEF double __clc_tgamma(double x) { return g; } -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, __clc_tgamma, double); +#define __DOUBLE_ONLY +#define FUNCTION __clc_tgamma +#define __CLC_BODY <clc/shared/unary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef FUNCTION #endif diff --git a/libclc/clc/lib/r600/math/clc_fmax.cl b/libclc/clc/lib/r600/math/clc_fmax.cl index 7e9c6df789eda..689e51a9829aa 100644 --- a/libclc/clc/lib/r600/math/clc_fmax.cl +++ b/libclc/clc/lib/r600/math/clc_fmax.cl @@ -17,7 +17,12 @@ _CLC_DEF _CLC_OVERLOAD float __clc_fmax(float x, float y) { y = __clc_flush_denormal_if_not_supported(y); return __builtin_fmaxf(x, y); } -_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_fmax, float, float) + +#define __FLOAT_ONLY +#define FUNCTION __clc_fmax +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef FUNCTION #ifdef cl_khr_fp64 @@ -26,7 +31,11 @@ _CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_fmax, float, float) _CLC_DEF _CLC_OVERLOAD double __clc_fmax(double x, double y) { return __builtin_fmax(x, y); } -_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, __clc_fmax, double, - double) + +#define __DOUBLE_ONLY +#define FUNCTION __clc_fmax +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef FUNCTION #endif diff --git a/libclc/clc/lib/r600/math/clc_fmin.cl b/libclc/clc/lib/r600/math/clc_fmin.cl index cbcf849f7b7e8..22cb7046a4ce3 100644 --- a/libclc/clc/lib/r600/math/clc_fmin.cl +++ b/libclc/clc/lib/r600/math/clc_fmin.cl @@ -18,7 +18,12 @@ _CLC_DEF _CLC_OVERLOAD float __clc_fmin(float x, float y) { y = __clc_flush_denormal_if_not_supported(y); return __builtin_fminf(x, y); } -_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_fmin, float, float) + +#define __FLOAT_ONLY +#define FUNCTION __clc_fmin +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef FUNCTION #ifdef cl_khr_fp64 @@ -27,7 +32,11 @@ _CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_fmin, float, float) _CLC_DEF _CLC_OVERLOAD double __clc_fmin(double x, double y) { return __builtin_fmin(x, y); } -_CLC_BINARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, __clc_fmin, double, - double) + +#define __DOUBLE_ONLY +#define FUNCTION __clc_fmin +#define __CLC_BODY <clc/shared/binary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef FUNCTION #endif diff --git a/libclc/clc/lib/r600/math/clc_native_rsqrt.cl b/libclc/clc/lib/r600/math/clc_native_rsqrt.cl index ee09814eb1e76..b5966570804c9 100644 --- a/libclc/clc/lib/r600/math/clc_native_rsqrt.cl +++ b/libclc/clc/lib/r600/math/clc_native_rsqrt.cl @@ -13,4 +13,7 @@ _CLC_OVERLOAD _CLC_DEF float __clc_native_rsqrt(float x) { return __builtin_r600_recipsqrt_ieeef(x); } -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_native_rsqrt, float); +#define __FLOAT_ONLY +#define FUNCTION __clc_native_rsqrt +#define __CLC_BODY <clc/shared/unary_def_scalarize.inc> +#include <clc/math/gentype.inc> diff --git a/libclc/clc/lib/r600/math/clc_rsqrt_override.cl b/libclc/clc/lib/r600/math/clc_rsqrt_override.cl index 91350d611f0ac..75355df56d32e 100644 --- a/libclc/clc/lib/r600/math/clc_rsqrt_override.cl +++ b/libclc/clc/lib/r600/math/clc_rsqrt_override.cl @@ -13,7 +13,11 @@ _CLC_OVERLOAD _CLC_DEF float __clc_rsqrt(float x) { return __builtin_r600_recipsqrt_ieeef(x); } -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, float, __clc_rsqrt, float); +#define __FLOAT_ONLY +#define FUNCTION __clc_rsqrt +#define __CLC_BODY <clc/shared/unary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef FUNCTION #ifdef cl_khr_fp64 @@ -23,6 +27,10 @@ _CLC_OVERLOAD _CLC_DEF double __clc_rsqrt(double x) { return __builtin_r600_recipsqrt_ieee(x); } -_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, double, __clc_rsqrt, double); +#define __DOUBLE_ONLY +#define FUNCTION __clc_rsqrt +#define __CLC_BODY <clc/shared/unary_def_scalarize.inc> +#include <clc/math/gentype.inc> +#undef FUNCTION #endif diff --git a/libclc/opencl/include/clc/opencl/geometric/fast_distance.h b/libclc/opencl/include/clc/opencl/geometric/fast_distance.h index ef401fdf7289f..975c485f16614 100644 --- a/libclc/opencl/include/clc/opencl/geometric/fast_distance.h +++ b/libclc/opencl/include/clc/opencl/geometric/fast_distance.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/geometric/fast_length.h b/libclc/opencl/include/clc/opencl/geometric/fast_length.h index b6fba6a8ca30c..1efa18ccd3061 100644 --- a/libclc/opencl/include/clc/opencl/geometric/fast_length.h +++ b/libclc/opencl/include/clc/opencl/geometric/fast_length.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/geometric/fast_normalize.h b/libclc/opencl/include/clc/opencl/geometric/fast_normalize.h index 1af81c66973b7..a860b9ceeb21d 100644 --- a/libclc/opencl/include/clc/opencl/geometric/fast_normalize.h +++ b/libclc/opencl/include/clc/opencl/geometric/fast_normalize.h @@ -15,4 +15,3 @@ #undef __CLC_FUNCTION #undef __CLC_GEOMETRIC_RET_GENTYPE -#undef __FLOAT_ONLY diff --git a/libclc/opencl/include/clc/opencl/math/half_cos.h b/libclc/opencl/include/clc/opencl/math/half_cos.h index c9a5a88f80910..f6d2966b2222a 100644 --- a/libclc/opencl/include/clc/opencl/math/half_cos.h +++ b/libclc/opencl/include/clc/opencl/math/half_cos.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/half_exp.h b/libclc/opencl/include/clc/opencl/math/half_exp.h index 65734878eb4e6..33d79c2e7d961 100644 --- a/libclc/opencl/include/clc/opencl/math/half_exp.h +++ b/libclc/opencl/include/clc/opencl/math/half_exp.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/half_exp10.h b/libclc/opencl/include/clc/opencl/math/half_exp10.h index 4df20a7aa6222..1c1a79e10c0de 100644 --- a/libclc/opencl/include/clc/opencl/math/half_exp10.h +++ b/libclc/opencl/include/clc/opencl/math/half_exp10.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/half_exp2.h b/libclc/opencl/include/clc/opencl/math/half_exp2.h index 2eed851fa698f..09c4c08986bff 100644 --- a/libclc/opencl/include/clc/opencl/math/half_exp2.h +++ b/libclc/opencl/include/clc/opencl/math/half_exp2.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/half_log.h b/libclc/opencl/include/clc/opencl/math/half_log.h index 3d44680570098..3bf0857ffad05 100644 --- a/libclc/opencl/include/clc/opencl/math/half_log.h +++ b/libclc/opencl/include/clc/opencl/math/half_log.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/half_log10.h b/libclc/opencl/include/clc/opencl/math/half_log10.h index 1c468f3f33985..117649055cad5 100644 --- a/libclc/opencl/include/clc/opencl/math/half_log10.h +++ b/libclc/opencl/include/clc/opencl/math/half_log10.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/half_log2.h b/libclc/opencl/include/clc/opencl/math/half_log2.h index 0695e2542962f..15f77f56d6177 100644 --- a/libclc/opencl/include/clc/opencl/math/half_log2.h +++ b/libclc/opencl/include/clc/opencl/math/half_log2.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/half_recip.h b/libclc/opencl/include/clc/opencl/math/half_recip.h index 8e2d16099f9f7..b1a48d4bbefbc 100644 --- a/libclc/opencl/include/clc/opencl/math/half_recip.h +++ b/libclc/opencl/include/clc/opencl/math/half_recip.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/half_rsqrt.h b/libclc/opencl/include/clc/opencl/math/half_rsqrt.h index e93973bc64575..4a277013df23a 100644 --- a/libclc/opencl/include/clc/opencl/math/half_rsqrt.h +++ b/libclc/opencl/include/clc/opencl/math/half_rsqrt.h @@ -10,5 +10,4 @@ #define __CLC_FUNCTION half_rsqrt #define __FLOAT_ONLY #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/half_sin.h b/libclc/opencl/include/clc/opencl/math/half_sin.h index 50626f9d96fe6..ad6309100ada0 100644 --- a/libclc/opencl/include/clc/opencl/math/half_sin.h +++ b/libclc/opencl/include/clc/opencl/math/half_sin.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/half_sqrt.h b/libclc/opencl/include/clc/opencl/math/half_sqrt.h index 7298570984baa..00d2036fb2d35 100644 --- a/libclc/opencl/include/clc/opencl/math/half_sqrt.h +++ b/libclc/opencl/include/clc/opencl/math/half_sqrt.h @@ -10,5 +10,4 @@ #define __CLC_FUNCTION half_sqrt #define __FLOAT_ONLY #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/half_tan.h b/libclc/opencl/include/clc/opencl/math/half_tan.h index 40ae6e6c9aba9..ee5f24ba55f5f 100644 --- a/libclc/opencl/include/clc/opencl/math/half_tan.h +++ b/libclc/opencl/include/clc/opencl/math/half_tan.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/native_cos.h b/libclc/opencl/include/clc/opencl/math/native_cos.h index e301634322b69..17967ac462c19 100644 --- a/libclc/opencl/include/clc/opencl/math/native_cos.h +++ b/libclc/opencl/include/clc/opencl/math/native_cos.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/native_exp.h b/libclc/opencl/include/clc/opencl/math/native_exp.h index bfe24aa12808a..51adca24b1bae 100644 --- a/libclc/opencl/include/clc/opencl/math/native_exp.h +++ b/libclc/opencl/include/clc/opencl/math/native_exp.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/native_exp10.h b/libclc/opencl/include/clc/opencl/math/native_exp10.h index f775478eef3a9..9c832735b0d97 100644 --- a/libclc/opencl/include/clc/opencl/math/native_exp10.h +++ b/libclc/opencl/include/clc/opencl/math/native_exp10.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/native_exp2.h b/libclc/opencl/include/clc/opencl/math/native_exp2.h index c00632c99387b..44bdbdee56adf 100644 --- a/libclc/opencl/include/clc/opencl/math/native_exp2.h +++ b/libclc/opencl/include/clc/opencl/math/native_exp2.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/native_log.h b/libclc/opencl/include/clc/opencl/math/native_log.h index 854f420723742..9ca7df63ecd9b 100644 --- a/libclc/opencl/include/clc/opencl/math/native_log.h +++ b/libclc/opencl/include/clc/opencl/math/native_log.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/native_log10.h b/libclc/opencl/include/clc/opencl/math/native_log10.h index 6862124b0f69a..10f9f6e87c5c1 100644 --- a/libclc/opencl/include/clc/opencl/math/native_log10.h +++ b/libclc/opencl/include/clc/opencl/math/native_log10.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/native_log2.h b/libclc/opencl/include/clc/opencl/math/native_log2.h index 576f50079d269..f481930cae61b 100644 --- a/libclc/opencl/include/clc/opencl/math/native_log2.h +++ b/libclc/opencl/include/clc/opencl/math/native_log2.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/native_recip.h b/libclc/opencl/include/clc/opencl/math/native_recip.h index 7bc99233fe008..6ab10e88cf712 100644 --- a/libclc/opencl/include/clc/opencl/math/native_recip.h +++ b/libclc/opencl/include/clc/opencl/math/native_recip.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/native_rsqrt.h b/libclc/opencl/include/clc/opencl/math/native_rsqrt.h index 40a1f02cba075..fc6e384e47998 100644 --- a/libclc/opencl/include/clc/opencl/math/native_rsqrt.h +++ b/libclc/opencl/include/clc/opencl/math/native_rsqrt.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/native_sin.h b/libclc/opencl/include/clc/opencl/math/native_sin.h index 798d5b325f6f1..295ac31dcf388 100644 --- a/libclc/opencl/include/clc/opencl/math/native_sin.h +++ b/libclc/opencl/include/clc/opencl/math/native_sin.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/native_sqrt.h b/libclc/opencl/include/clc/opencl/math/native_sqrt.h index a41a0c104cc84..76beaa3814e00 100644 --- a/libclc/opencl/include/clc/opencl/math/native_sqrt.h +++ b/libclc/opencl/include/clc/opencl/math/native_sqrt.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION diff --git a/libclc/opencl/include/clc/opencl/math/native_tan.h b/libclc/opencl/include/clc/opencl/math/native_tan.h index 658bb8f932743..02afc2b4a1316 100644 --- a/libclc/opencl/include/clc/opencl/math/native_tan.h +++ b/libclc/opencl/include/clc/opencl/math/native_tan.h @@ -12,5 +12,4 @@ #include <clc/math/gentype.inc> -#undef __FLOAT_ONLY #undef __CLC_FUNCTION _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
