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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
