[libclc] [NFC][libclc] Refactor _CLC_*_VECTORIZE macros to functions in .inc files (PR #145678)
Wenju He via cfe-commits
cfe-commits at lists.llvm.org
Wed Jun 25 04:26:40 PDT 2025
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.he at 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
More information about the cfe-commits
mailing list