[libclc] [libclc] Move clcmacro.h to CLC library. NFC (PR #114845)
via cfe-commits
cfe-commits at lists.llvm.org
Mon Nov 4 10:22:06 PST 2024
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Fraser Cormack (frasercrmck)
<details>
<summary>Changes</summary>
---
Patch is 50.86 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/114845.diff
95 Files Affected:
- (modified) libclc/amdgcn/lib/integer/popcount.cl (+1-1)
- (modified) libclc/amdgcn/lib/math/fmax.cl (+1-2)
- (modified) libclc/amdgcn/lib/math/fmin.cl (+1-2)
- (modified) libclc/amdgcn/lib/math/ldexp.cl (+1-2)
- (modified) libclc/amdgpu/lib/math/half_native_unary.inc (+1-1)
- (modified) libclc/amdgpu/lib/math/nextafter.cl (+1-1)
- (modified) libclc/amdgpu/lib/math/sqrt.cl (+2-2)
- (renamed) libclc/clc/include/clc/clcmacro.h (+116-98)
- (modified) libclc/clspv/lib/math/fma.cl (+2-2)
- (modified) libclc/generic/include/math/clc_sqrt.h (+3)
- (removed) libclc/generic/include/utils.h (-10)
- (modified) libclc/generic/lib/atom_int32_binary.inc (+1-1)
- (modified) libclc/generic/lib/common/degrees.cl (+1-2)
- (modified) libclc/generic/lib/common/radians.cl (+1-2)
- (modified) libclc/generic/lib/common/sign.cl (+1-1)
- (modified) libclc/generic/lib/common/smoothstep.cl (+1-2)
- (modified) libclc/generic/lib/common/step.cl (+1-2)
- (modified) libclc/generic/lib/integer/add_sat.cl (+1-1)
- (modified) libclc/generic/lib/integer/clz.cl (+1-1)
- (modified) libclc/generic/lib/integer/mad_sat.cl (+1-1)
- (modified) libclc/generic/lib/integer/sub_sat.cl (+1-1)
- (modified) libclc/generic/lib/math/acos.cl (+1-1)
- (modified) libclc/generic/lib/math/acosh.cl (+1-1)
- (modified) libclc/generic/lib/math/acospi.cl (+1-1)
- (modified) libclc/generic/lib/math/asin.cl (+1-1)
- (modified) libclc/generic/lib/math/asinh.cl (+1-1)
- (modified) libclc/generic/lib/math/asinpi.cl (+1-1)
- (modified) libclc/generic/lib/math/atan.cl (+1-1)
- (modified) libclc/generic/lib/math/atan2.cl (+1-1)
- (modified) libclc/generic/lib/math/atan2pi.cl (+1-1)
- (modified) libclc/generic/lib/math/atanh.cl (+1-1)
- (modified) libclc/generic/lib/math/atanpi.cl (+1-1)
- (modified) libclc/generic/lib/math/cbrt.cl (+1-1)
- (modified) libclc/generic/lib/math/ceil.cl (+1-1)
- (modified) libclc/generic/lib/math/clc_exp10.cl (+1-1)
- (modified) libclc/generic/lib/math/clc_fma.cl (+1-1)
- (modified) libclc/generic/lib/math/clc_fmod.cl (+1-1)
- (modified) libclc/generic/lib/math/clc_hypot.cl (+1-1)
- (modified) libclc/generic/lib/math/clc_ldexp.cl (+1-1)
- (modified) libclc/generic/lib/math/clc_nextafter.cl (+1-1)
- (modified) libclc/generic/lib/math/clc_pow.cl (+1-1)
- (modified) libclc/generic/lib/math/clc_pown.cl (+1-1)
- (modified) libclc/generic/lib/math/clc_powr.cl (+1-1)
- (modified) libclc/generic/lib/math/clc_remainder.cl (+1-1)
- (modified) libclc/generic/lib/math/clc_remquo.cl (+1-1)
- (modified) libclc/generic/lib/math/clc_rootn.cl (+1-1)
- (modified) libclc/generic/lib/math/clc_sw_binary.inc (+1-1)
- (modified) libclc/generic/lib/math/clc_sw_unary.inc (+1-1)
- (modified) libclc/generic/lib/math/clc_tan.cl (+1-1)
- (modified) libclc/generic/lib/math/clc_tanpi.cl (+1-1)
- (modified) libclc/generic/lib/math/copysign.cl (+1-1)
- (modified) libclc/generic/lib/math/cos.cl (+1-1)
- (modified) libclc/generic/lib/math/cosh.cl (+1-1)
- (modified) libclc/generic/lib/math/cospi.cl (+1-1)
- (modified) libclc/generic/lib/math/erf.cl (+1-1)
- (modified) libclc/generic/lib/math/erfc.cl (+1-1)
- (modified) libclc/generic/lib/math/exp.cl (+1-1)
- (modified) libclc/generic/lib/math/exp2.cl (+1-1)
- (modified) libclc/generic/lib/math/expm1.cl (+1-1)
- (modified) libclc/generic/lib/math/fabs.cl (+1-1)
- (modified) libclc/generic/lib/math/floor.cl (+1-1)
- (modified) libclc/generic/lib/math/fmax.cl (+1-2)
- (modified) libclc/generic/lib/math/fmin.cl (+1-2)
- (modified) libclc/generic/lib/math/frexp.cl (+1-1)
- (modified) libclc/generic/lib/math/frexp.inc (+1-1)
- (modified) libclc/generic/lib/math/half_binary.inc (+1-1)
- (modified) libclc/generic/lib/math/half_unary.inc (+1-1)
- (modified) libclc/generic/lib/math/ilogb.cl (+2-2)
- (modified) libclc/generic/lib/math/ldexp.cl (+2-2)
- (modified) libclc/generic/lib/math/lgamma.cl (+1-1)
- (modified) libclc/generic/lib/math/lgamma_r.cl (+1-1)
- (modified) libclc/generic/lib/math/log.cl (+1-1)
- (modified) libclc/generic/lib/math/log10.cl (+2-2)
- (modified) libclc/generic/lib/math/log1p.cl (+1-1)
- (modified) libclc/generic/lib/math/log2.cl (+2-2)
- (modified) libclc/generic/lib/math/logb.cl (+2-2)
- (modified) libclc/generic/lib/math/maxmag.cl (+1-1)
- (modified) libclc/generic/lib/math/minmag.cl (+1-1)
- (modified) libclc/generic/lib/math/nan.cl (+1-1)
- (modified) libclc/generic/lib/math/native_unary_intrinsic.inc (+1-1)
- (modified) libclc/generic/lib/math/rsqrt.cl (+1-2)
- (modified) libclc/generic/lib/math/sin.cl (+1-1)
- (modified) libclc/generic/lib/math/sinh.cl (+1-1)
- (modified) libclc/generic/lib/math/sinpi.cl (+1-1)
- (modified) libclc/generic/lib/math/tables.h (+2)
- (modified) libclc/generic/lib/math/tanh.cl (+1-1)
- (modified) libclc/generic/lib/math/tgamma.cl (+1-1)
- (modified) libclc/generic/lib/math/unary_builtin.inc (+2-2)
- (modified) libclc/generic/lib/relational/bitselect.cl (+1-2)
- (modified) libclc/generic/lib/relational/select.cl (+1-1)
- (modified) libclc/ptx/lib/math/nextafter.cl (+1-1)
- (modified) libclc/r600/lib/math/fmax.cl (+1-1)
- (modified) libclc/r600/lib/math/fmin.cl (+1-1)
- (modified) libclc/r600/lib/math/native_rsqrt.cl (+1-2)
- (modified) libclc/r600/lib/math/rsqrt.cl (+1-2)
``````````diff
diff --git a/libclc/amdgcn/lib/integer/popcount.cl b/libclc/amdgcn/lib/integer/popcount.cl
index ebd167d04da7cb..3b493fbd146f01 100644
--- a/libclc/amdgcn/lib/integer/popcount.cl
+++ b/libclc/amdgcn/lib/integer/popcount.cl
@@ -1,5 +1,5 @@
#include <clc/clc.h>
-#include <utils.h>
+#include <clc/utils.h>
#include <integer/popcount.h>
#define __CLC_BODY "popcount.inc"
diff --git a/libclc/amdgcn/lib/math/fmax.cl b/libclc/amdgcn/lib/math/fmax.cl
index cb796161010829..4407d4a87f9ea9 100644
--- a/libclc/amdgcn/lib/math/fmax.cl
+++ b/libclc/amdgcn/lib/math/fmax.cl
@@ -1,6 +1,5 @@
#include <clc/clc.h>
-
-#include "../../../generic/lib/clcmacro.h"
+#include <clc/clcmacro.h>
_CLC_DEF _CLC_OVERLOAD float fmax(float x, float y)
{
diff --git a/libclc/amdgcn/lib/math/fmin.cl b/libclc/amdgcn/lib/math/fmin.cl
index 35dea8bc975f7f..4d02a47babdabb 100644
--- a/libclc/amdgcn/lib/math/fmin.cl
+++ b/libclc/amdgcn/lib/math/fmin.cl
@@ -1,6 +1,5 @@
#include <clc/clc.h>
-
-#include "../../../generic/lib/clcmacro.h"
+#include <clc/clcmacro.h>
_CLC_DEF _CLC_OVERLOAD float fmin(float x, float y)
{
diff --git a/libclc/amdgcn/lib/math/ldexp.cl b/libclc/amdgcn/lib/math/ldexp.cl
index 9713e4daee0e79..d46d2dc606818a 100644
--- a/libclc/amdgcn/lib/math/ldexp.cl
+++ b/libclc/amdgcn/lib/math/ldexp.cl
@@ -21,8 +21,7 @@
*/
#include <clc/clc.h>
-
-#include "../../../generic/lib/clcmacro.h"
+#include <clc/clcmacro.h>
#ifdef __HAS_LDEXPF__
#define BUILTINF __builtin_amdgcn_ldexpf
diff --git a/libclc/amdgpu/lib/math/half_native_unary.inc b/libclc/amdgpu/lib/math/half_native_unary.inc
index 0f99ba5e521630..bdc380600501a6 100644
--- a/libclc/amdgpu/lib/math/half_native_unary.inc
+++ b/libclc/amdgpu/lib/math/half_native_unary.inc
@@ -1,4 +1,4 @@
-#include <utils.h>
+#include <clc/utils.h>
#define __CLC_HALF_FUNC(x) __CLC_CONCAT(half_, x)
#define __CLC_NATIVE_FUNC(x) __CLC_CONCAT(native_, x)
diff --git a/libclc/amdgpu/lib/math/nextafter.cl b/libclc/amdgpu/lib/math/nextafter.cl
index b290da0e417dd4..6dc117b8cdd64c 100644
--- a/libclc/amdgpu/lib/math/nextafter.cl
+++ b/libclc/amdgpu/lib/math/nextafter.cl
@@ -1,5 +1,5 @@
#include <clc/clc.h>
-#include "../lib/clcmacro.h"
+#include <clc/clcmacro.h>
#include <math/clc_nextafter.h>
_CLC_DEFINE_BINARY_BUILTIN(float, nextafter, __clc_nextafter, float, float)
diff --git a/libclc/amdgpu/lib/math/sqrt.cl b/libclc/amdgpu/lib/math/sqrt.cl
index 556260033169b6..17d77e50d44d3f 100644
--- a/libclc/amdgpu/lib/math/sqrt.cl
+++ b/libclc/amdgpu/lib/math/sqrt.cl
@@ -20,9 +20,9 @@
* THE SOFTWARE.
*/
-#include <clc/clc.h>
-#include "../../../generic/lib/clcmacro.h"
#include "math/clc_sqrt.h"
+#include <clc/clc.h>
+#include <clc/clcmacro.h>
_CLC_DEFINE_UNARY_BUILTIN(float, sqrt, __clc_sqrt, float)
diff --git a/libclc/generic/lib/clcmacro.h b/libclc/clc/include/clc/clcmacro.h
similarity index 55%
rename from libclc/generic/lib/clcmacro.h
rename to libclc/clc/include/clc/clcmacro.h
index fa302b398adc68..244239284ecabc 100644
--- a/libclc/generic/lib/clcmacro.h
+++ b/libclc/clc/include/clc/clcmacro.h
@@ -1,92 +1,105 @@
-#include <clc/clc.h>
-#include <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.x), FUNCTION(x.y)); \
- } \
-\
- DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x) { \
- return (RET_TYPE##3)(FUNCTION(x.x), FUNCTION(x.y), FUNCTION(x.z)); \
- } \
-\
- DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x) { \
- return (RET_TYPE##4)(FUNCTION(x.lo), FUNCTION(x.hi)); \
- } \
-\
- DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x) { \
- return (RET_TYPE##8)(FUNCTION(x.lo), FUNCTION(x.hi)); \
- } \
-\
- DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x) { \
- return (RET_TYPE##16)(FUNCTION(x.lo), FUNCTION(x.hi)); \
+#ifndef __CLC_CLCMACRO_H__
+#define __CLC_CLCMACRO_H__
+
+#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.x), FUNCTION(x.y)); \
+ } \
+ \
+ DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x) { \
+ return (RET_TYPE##3)(FUNCTION(x.x), FUNCTION(x.y), FUNCTION(x.z)); \
+ } \
+ \
+ DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x) { \
+ return (RET_TYPE##4)(FUNCTION(x.lo), FUNCTION(x.hi)); \
+ } \
+ \
+ DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x) { \
+ return (RET_TYPE##8)(FUNCTION(x.lo), FUNCTION(x.hi)); \
+ } \
+ \
+ DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x) { \
+ return (RET_TYPE##16)(FUNCTION(x.lo), FUNCTION(x.hi)); \
}
-#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.x, y.x), FUNCTION(x.y, y.y)); \
- } \
-\
- DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x, ARG2_TYPE##3 y) { \
- return (RET_TYPE##3)(FUNCTION(x.x, y.x), FUNCTION(x.y, y.y), \
- FUNCTION(x.z, y.z)); \
- } \
-\
- DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x, ARG2_TYPE##4 y) { \
- return (RET_TYPE##4)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
- } \
-\
- DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x, ARG2_TYPE##8 y) { \
- return (RET_TYPE##8)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
- } \
-\
- DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x, ARG2_TYPE##16 y) { \
- return (RET_TYPE##16)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
+#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.x, y.x), FUNCTION(x.y, y.y)); \
+ } \
+ \
+ DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x, ARG2_TYPE##3 y) { \
+ return (RET_TYPE##3)(FUNCTION(x.x, y.x), FUNCTION(x.y, y.y), \
+ FUNCTION(x.z, y.z)); \
+ } \
+ \
+ DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x, ARG2_TYPE##4 y) { \
+ return (RET_TYPE##4)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
+ } \
+ \
+ DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x, ARG2_TYPE##8 y) { \
+ return (RET_TYPE##8)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
+ } \
+ \
+ DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x, ARG2_TYPE##16 y) { \
+ return (RET_TYPE##16)(FUNCTION(x.lo, y.lo), FUNCTION(x.hi, y.hi)); \
}
-#define _CLC_V_S_V_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, ARG2_TYPE) \
- DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE x, ARG2_TYPE##2 y) { \
- return (RET_TYPE##2)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
- } \
-\
- DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE x, ARG2_TYPE##3 y) { \
- return (RET_TYPE##3)(FUNCTION(x, y.x), FUNCTION(x, y.y), \
- FUNCTION(x, y.z)); \
- } \
-\
- DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE x, ARG2_TYPE##4 y) { \
- return (RET_TYPE##4)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
- } \
-\
- DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE x, ARG2_TYPE##8 y) { \
- return (RET_TYPE##8)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
- } \
-\
- DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE x, ARG2_TYPE##16 y) { \
- return (RET_TYPE##16)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
- } \
-\
-
-#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.x, y.x, z.x), FUNCTION(x.y, y.y, z.y)); \
- } \
-\
- DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x, ARG2_TYPE##3 y, ARG3_TYPE##3 z) { \
- return (RET_TYPE##3)(FUNCTION(x.x, y.x, z.x), FUNCTION(x.y, y.y, z.y), \
- FUNCTION(x.z, y.z, z.z)); \
- } \
-\
- DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x, ARG2_TYPE##4 y, ARG3_TYPE##4 z) { \
- return (RET_TYPE##4)(FUNCTION(x.lo, y.lo, z.lo), FUNCTION(x.hi, y.hi, z.hi)); \
- } \
-\
- DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x, ARG2_TYPE##8 y, ARG3_TYPE##8 z) { \
- return (RET_TYPE##8)(FUNCTION(x.lo, y.lo, z.lo), FUNCTION(x.hi, y.hi, z.hi)); \
- } \
-\
- DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x, ARG2_TYPE##16 y, ARG3_TYPE##16 z) { \
- return (RET_TYPE##16)(FUNCTION(x.lo, y.lo, z.lo), FUNCTION(x.hi, y.hi, z.hi)); \
+#define _CLC_V_S_V_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, \
+ ARG2_TYPE) \
+ DECLSPEC RET_TYPE##2 FUNCTION(ARG1_TYPE x, ARG2_TYPE##2 y) { \
+ return (RET_TYPE##2)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
+ } \
+ \
+ DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE x, ARG2_TYPE##3 y) { \
+ return (RET_TYPE##3)(FUNCTION(x, y.x), FUNCTION(x, y.y), \
+ FUNCTION(x, y.z)); \
+ } \
+ \
+ DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE x, ARG2_TYPE##4 y) { \
+ return (RET_TYPE##4)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
+ } \
+ \
+ DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE x, ARG2_TYPE##8 y) { \
+ return (RET_TYPE##8)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
+ } \
+ \
+ DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE x, ARG2_TYPE##16 y) { \
+ return (RET_TYPE##16)(FUNCTION(x, y.lo), FUNCTION(x, y.hi)); \
+ }
+
+#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.x, y.x, z.x), FUNCTION(x.y, y.y, z.y)); \
+ } \
+ \
+ DECLSPEC RET_TYPE##3 FUNCTION(ARG1_TYPE##3 x, ARG2_TYPE##3 y, \
+ ARG3_TYPE##3 z) { \
+ return (RET_TYPE##3)(FUNCTION(x.x, y.x, z.x), FUNCTION(x.y, y.y, z.y), \
+ FUNCTION(x.z, y.z, z.z)); \
+ } \
+ \
+ DECLSPEC RET_TYPE##4 FUNCTION(ARG1_TYPE##4 x, ARG2_TYPE##4 y, \
+ ARG3_TYPE##4 z) { \
+ return (RET_TYPE##4)(FUNCTION(x.lo, y.lo, z.lo), \
+ FUNCTION(x.hi, y.hi, z.hi)); \
+ } \
+ \
+ DECLSPEC RET_TYPE##8 FUNCTION(ARG1_TYPE##8 x, ARG2_TYPE##8 y, \
+ ARG3_TYPE##8 z) { \
+ return (RET_TYPE##8)(FUNCTION(x.lo, y.lo, z.lo), \
+ FUNCTION(x.hi, y.hi, z.hi)); \
+ } \
+ \
+ DECLSPEC RET_TYPE##16 FUNCTION(ARG1_TYPE##16 x, ARG2_TYPE##16 y, \
+ ARG3_TYPE##16 z) { \
+ return (RET_TYPE##16)(FUNCTION(x.lo, y.lo, z.lo), \
+ FUNCTION(x.hi, y.hi, z.hi)); \
}
#define _CLC_V_S_S_V_VECTORIZE(DECLSPEC, RET_TYPE, FUNCTION, ARG1_TYPE, \
@@ -161,21 +174,24 @@
ARG2_TYPE, 8) *)((ADDR_SPACE ARG2_TYPE *)y + 8))); \
}
-#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)
+#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)
-#define _CLC_DEFINE_BINARY_BUILTIN_WITH_SCALAR_SECOND_ARG(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, ARG2_TYPE) \
-_CLC_DEFINE_BINARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, ARG2_TYPE) \
-_CLC_BINARY_VECTORIZE_SCALAR_SECOND_ARG(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE, ARG2_TYPE)
+#define _CLC_DEFINE_BINARY_BUILTIN_WITH_SCALAR_SECOND_ARG( \
+ RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, ARG2_TYPE) \
+ _CLC_DEFINE_BINARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE, \
+ ARG2_TYPE) \
+ _CLC_BINARY_VECTORIZE_SCALAR_SECOND_ARG(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, \
+ FUNCTION, ARG1_TYPE, ARG2_TYPE)
-#define _CLC_DEFINE_UNARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE) \
-_CLC_DEF _CLC_OVERLOAD RET_TYPE FUNCTION(ARG1_TYPE x) { \
- return BUILTIN(x); \
-} \
-_CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE)
+#define _CLC_DEFINE_UNARY_BUILTIN(RET_TYPE, FUNCTION, BUILTIN, ARG1_TYPE) \
+ _CLC_DEF _CLC_OVERLOAD RET_TYPE FUNCTION(ARG1_TYPE x) { return BUILTIN(x); } \
+ _CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE)
#ifdef cl_khr_fp16
@@ -199,3 +215,5 @@ _CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, RET_TYPE, FUNCTION, ARG1_TYPE)
#define _CLC_DEFINE_BINARY_BUILTIN_FP16(FUNCTION)
#endif
+
+#endif // __CLC_CLCMACRO_H__
diff --git a/libclc/clspv/lib/math/fma.cl b/libclc/clspv/lib/math/fma.cl
index 3ffca28bd3bef6..e6251db4e92dbe 100644
--- a/libclc/clspv/lib/math/fma.cl
+++ b/libclc/clspv/lib/math/fma.cl
@@ -24,9 +24,9 @@
// (__clc_sw_fma), but avoids the use of ulong in favor of uint2. The logic has
// been updated as appropriate.
-#include <clc/clc.h>
-#include "../../../generic/lib/clcmacro.h"
#include "../../../generic/lib/math/math.h"
+#include <clc/clc.h>
+#include <clc/clcmacro.h>
struct fp {
uint2 mantissa;
diff --git a/libclc/generic/include/math/clc_sqrt.h b/libclc/generic/include/math/clc_sqrt.h
index 60e183ff66a5ce..90a7c575c9bad9 100644
--- a/libclc/generic/include/math/clc_sqrt.h
+++ b/libclc/generic/include/math/clc_sqrt.h
@@ -1,3 +1,6 @@
+#include <clc/clcfunc.h>
+#include <clc/clctypes.h>
+
#define __CLC_FUNCTION __clc_sqrt
#define __CLC_BODY <clc/math/unary_decl.inc>
#include <clc/math/gentype.inc>
diff --git a/libclc/generic/include/utils.h b/libclc/generic/include/utils.h
deleted file mode 100644
index 018a7b31f8ff3d..00000000000000
--- a/libclc/generic/include/utils.h
+++ /dev/null
@@ -1,10 +0,0 @@
-#ifndef __CLC_UTILS_H_
-#define __CLC_UTILS_H_
-
-#define __CLC_CONCAT(x, y) x ## y
-#define __CLC_XCONCAT(x, y) __CLC_CONCAT(x, y)
-
-#define __CLC_STR(x) #x
-#define __CLC_XSTR(x) __CLC_STR(x)
-
-#endif
diff --git a/libclc/generic/lib/atom_int32_binary.inc b/libclc/generic/lib/atom_int32_binary.inc
index 3af4c4bb0ae80f..5d3b33fcc1edfd 100644
--- a/libclc/generic/lib/atom_int32_binary.inc
+++ b/libclc/generic/lib/atom_int32_binary.inc
@@ -1,5 +1,5 @@
#include <clc/clc.h>
-#include "utils.h"
+#include <clc/utils.h>
#define __CLC_ATOM_IMPL(AS, TYPE) \
_CLC_OVERLOAD _CLC_DEF TYPE __CLC_XCONCAT(atom_, __CLC_ATOMIC_OP) (volatile AS TYPE *p, TYPE val) { \
diff --git a/libclc/generic/lib/common/degrees.cl b/libclc/generic/lib/common/degrees.cl
index 5de56f86c4ca93..cf49b190c76b37 100644
--- a/libclc/generic/lib/common/degrees.cl
+++ b/libclc/generic/lib/common/degrees.cl
@@ -21,8 +21,7 @@
*/
#include <clc/clc.h>
-
-#include "../clcmacro.h"
+#include <clc/clcmacro.h>
_CLC_OVERLOAD _CLC_DEF float degrees(float radians) {
// 180/pi = ~57.29577951308232087685 or 0x1.ca5dc1a63c1f8p+5 or 0x1.ca5dc2p+5F
diff --git a/libclc/generic/lib/common/radians.cl b/libclc/generic/lib/common/radians.cl
index 3838dd6cde60f0..645a30549afedd 100644
--- a/libclc/generic/lib/common/radians.cl
+++ b/libclc/generic/lib/common/radians.cl
@@ -21,8 +21,7 @@
*/
#include <clc/clc.h>
-
-#include "../clcmacro.h"
+#include <clc/clcmacro.h>
_CLC_OVERLOAD _CLC_DEF float radians(float degrees) {
// pi/180 = ~0.01745329251994329577 or 0x1.1df46a2529d39p-6 or 0x1.1df46ap-6F
diff --git a/libclc/generic/lib/common/sign.cl b/libclc/generic/lib/common/sign.cl
index 7b6b793be79c19..ad8f7405e0cb30 100644
--- a/libclc/generic/lib/common/sign.cl
+++ b/libclc/generic/lib/common/sign.cl
@@ -1,5 +1,5 @@
#include <clc/clc.h>
-#include "../clcmacro.h"
+#include <clc/clcmacro.h>
#define SIGN(TYPE, F) \
_CLC_DEF _CLC_OVERLOAD TYPE sign(TYPE x) { \
diff --git a/libclc/generic/lib/common/smoothstep.cl b/libclc/generic/lib/common/smoothstep.cl
index 1b6a74b89d2c21..4cdecfc4abe26e 100644
--- a/libclc/generic/lib/common/smoothstep.cl
+++ b/libclc/generic/lib/common/smoothstep.cl
@@ -21,8 +21,7 @@
*/
#include <clc/clc.h>
-
-#include "../clcmacro.h"
+#include <clc/clcmacro.h>
_CLC_OVERLOAD _CLC_DEF float smoothstep(float edge0, float edge1, float x) {
float t = clamp((x - edge0) / (edge1 - edge0), 0.0f, 1.0f);
diff --git a/libclc/generic/lib/common/step.cl b/libclc/generic/lib/common/step.cl
index 8155b469fb210f..3d9bc53c2e146d 100644
--- a/libclc/generic/lib/common/step.cl
+++ b/libclc/generic/lib/common/step.cl
@@ -21,8 +21,7 @@
*/
#include <clc/clc.h>
-
-#include "../clcmacro.h"
+#include <clc/clcmacr...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/114845
More information about the cfe-commits
mailing list