[libclc] Move clcmacro.h to CLC library. NFC #114845
Merged
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
Sorry for the noise, but this felt too big to just push without review. It should hopefully minimize diffs in upcoming patches. |
|
@llvm/pr-subscribers-backend-amdgpu Author: Fraser Cormack (frasercrmck) ChangesPatch is 50.86 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/114845.diff 95 Files Affected:
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]
|
arsenm
approved these changes
Nov 4, 2024
PhilippRados
pushed a commit
to PhilippRados/llvm-project
that referenced
this pull request
Nov 6, 2024
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
No description provided.